o
    (ºiÚ%  ã                   @   sú   d dl Z d dlmZ dd„ Zdd„ Ze jdd„ ƒZe jeƒ g d¢d	e jed
dej	dej	dej	dej	dej	f
dd„ƒƒZ
e jed
dej	dej	dej	dej	dej	dej	fdd„ƒZe jeƒ g d¢d	e jed
dej	dej	dej	dej	fdd„ƒƒZdS )é    Nc                   C   s   dd„ dD ƒS )Nc                 S   sN   g | ]#}d D ]}dD ]}dD ]}dD ]}t j|||ddœ||d‘qqq
qqS )©é€   )é@   )é   )é   é   )ÚBLOCK_SIZE_MÚBLOCK_SIZE_NÚBLOCK_SIZE_KÚGROUP_SIZE_M)Ú
num_stagesÚ	num_warps)ÚtritonÚConfig)Ú.0ÚBMÚBNÚBKÚsÚw© r   úi/home/ubuntu/veenaModal/venv/lib/python3.10/site-packages/flashinfer/triton/kernels/sm_constraint_gemm.pyÚ
<listcomp>   s0    ôóòñòüøÿz&matmul_get_configs.<locals>.<listcomp>r   r   r   r   r   r   Úmatmul_get_configs   s   õr   c                 C   sª   i }|d |d |d }}}|j › d|› d|› d|› d|d< d	|v r,|d	  ¡ }n|d
 r2dnd}d| | | |d|d › < ||| ||  ||   |d< |S )NÚMÚNÚKz [M=z, N=z, K=ú]ÚnameÚc_ptrÚ
FP8_OUTPUTé   é   g       @Úflopsr   Úbytes)r   Úelement_size)ÚgridÚkernelÚargsÚretr   r   r   Úbytes_per_elemr   r   r   Ú_matmul_launch_metadata   s   " r+   c           
      C   s>   | | }|| }t || |ƒ}|| |  }| | | }	||	fS )N)Úmin)
Útile_idÚnum_pid_in_groupÚ	num_pid_mr   ÚNUM_SMSÚgroup_idÚfirst_pid_mÚgroup_size_mÚpid_mÚpid_nr   r   r   Ú_compute_pid&   s   r6   )r   r   r   )ÚconfigsÚkey)Úlaunch_metadatar   r	   r
   r   r0   c           .   
   C   sÀ  t jdd}t  ||¡}t  ||¡}t  ||¡}|| }|| }t  d|¡}|| }t j|||ddD ])}t|||||ƒ\}}|| }|| }|t  d|¡ } |t  d|¡ }!t  | |k | d¡} t  |!|k |!d¡}!t  t  | |¡|¡} t  t  |!|¡|¡}!t j	||ft j
d}"t|ƒD ]g}#|#| t  d|¡ }$| | d d …d f | |$d d d …f |   }%||$d d …d f | |!d d d …f |	   }&t j|%|d d d …f ||#|  k dd}'t j|&|d d …d f ||#|  k dd}(t  |'|(|"¡}"qŒ||7 }t|||||ƒ\}}|| t  d|¡ })|| t  d|¡ }*||
|)d d …d f   ||*d d d …f   }+|)d d …d f |k |*d d d …f |k @ },|" |jj¡}-t  |-||t j|+|,d ¡}-t j|+|-|,d q3d S )	Nr   ©ÚaxisT©Úflatten©Údtypeç        ©ÚmaskÚother©rB   )ÚtlÚ
program_idÚcdivÚarangeÚranger6   ÚwhereÚmax_contiguousÚmultiple_ofÚzerosÚfloat32ÚloadÚdotÚtor?   Ú
element_tyÚfmaÚstore).Úa_ptrÚb_ptrr   r   r   r   Ú	stride_amÚ	stride_akÚ	stride_bkÚ	stride_bnÚ	stride_cmÚ	stride_cnÚalphaÚbetar   r	   r
   r   r0   Ú	start_pidr/   Ú	num_pid_nÚk_tilesÚ	num_tilesÚ	tile_id_cÚoffs_k_for_maskr.   r-   r4   r5   Ústart_mÚstart_nÚoffs_amÚoffs_bnÚaccumulatorÚkiÚoffs_kÚa_ptrsÚb_ptrsÚaÚbÚoffs_cmÚoffs_cnÚc_ptrsÚc_maskÚcr   r   r   Úgemm_kernel_persistent0   s`   
ÿ&ÿ&ÿÿÿ
ÿ,(Öru   ÚEPILOGUE_SUBTILEc           +      C   sb  |j j}tjdd}t ||¡}t ||	¡}t ||
¡}|| }tj| ||g|dg||
gd}tj|||g|dg|	|
gd}tj|||g|dg||sL|	n|	d gd}|| }|| }tj|||ddD ]Ê}t|||||ƒ\}}|| }||	 }tj||	ftj	d}t|ƒD ]}||
 } | 
|| g¡}!| 
|| g¡}"t |!|"j|¡}q†||7 }t|||||ƒ\}}|| }#||	 }$|rt ||d|	d f¡}%t |%d	¡}%t |%¡\}&}'t |&||| 
|#|$g¡ ¡}&t |'||| 
|#|$|	d  g¡ ¡}'|& |¡}(| |#|$g|(¡ |' |¡})| |#|$|	d  g|)¡ qdt |||| 
|#|$g¡ ¡}| |¡}*| |#|$g|*¡ qdd S )
Nr   r:   r!   )ÚshapeÚstridesÚblock_shaper"   Tr<   r>   )r   r"   r!   )r?   rR   rE   rF   rG   Úmake_tensor_descriptorrI   r6   rM   rN   rO   rP   ÚTÚreshapeÚpermuteÚsplitrS   rQ   rT   )+rU   rV   r   r   r   r   r]   r^   r   r	   r
   r   rv   r0   r?   r_   r/   r`   ra   rb   Úa_descÚb_descÚc_descrc   r.   r-   r4   r5   rg   rh   ri   rj   rk   rn   ro   rp   rq   ÚaccÚacc0Úacc1Úc0Úc1rt   r   r   r   Ú!gemm_kernel_descriptor_persistent„   s€   üüþü
ÿ
ÿÿ

ÿ
Úr‡   c           +   	   C   sž  t jdd}t  ||¡}t  ||¡}|| }|| }|| }t|| |ƒ}|||  }|| | }|| }|| }|t  d|¡ }|t  d|¡ }t  ||k |d¡}t  ||k |d¡}t  t  ||¡|¡}t  t  ||¡|¡}t  d|¡}| |d d …d f | |d d d …f |   } ||d d …d f | |d d d …f |	   }!t j||ft j	d}"t
dt  ||¡ƒD ]=}#t j| |d d d …f ||#|  k dd}$t j|!|d d …d f ||#|  k dd}%t  |$|%|"¡}"| || 7 } |!|| 7 }!q´|" |jj¡}&|| t  d|¡ }'|| t  d|¡ }(||
|'d d …d f   ||(d d d …f   })|'d d …d f |k |(d d d …f |k @ }*t  |&||t j|)|*d ¡}&t j|)|&|*d d S )Nr   r:   r>   r@   rA   rD   )rE   rF   rG   r,   rH   rJ   rK   rL   rM   rN   rI   rO   rP   rQ   r?   rR   rS   rT   )+rU   rV   r   r   r   r   rW   rX   rY   rZ   r[   r\   r]   r^   r   r	   r
   r   Úpidr/   r`   r.   r1   r2   r3   r4   r5   re   rf   rg   rh   rk   rl   rm   ri   Úkrn   ro   rt   rp   rq   rr   rs   r   r   r   Úgemm_kernelá   sD   ,,((,(rŠ   )r   Útriton.languageÚlanguagerE   r   r+   Újitr6   ÚautotuneÚ	constexprru   r‡   rŠ   r   r   r   r   Ú<module>   sd    
	þ
ñðïîí
O	÷
öõôóò\þ
ñðïî