
    )`i%                     B   d dl Z d dlmZ d Zd Ze j        d             Z e j         e            g d           e j        e          dej	        d	ej	        d
ej	        dej	        dej	        f
d                        Z
 e j        e          dej	        d	ej	        d
ej	        dej	        dej	        dej	        fd            Z e j         e            g d           e j        e          dej	        d	ej	        d
ej	        dej	        fd                        ZdS )    Nc                      d dD             S )Nc                 n    g | ]2}d D ]-}dD ](}dD ]#}dD ]}t          j        |||dd||          $).3S )   )@   )   )      )BLOCK_SIZE_MBLOCK_SIZE_NBLOCK_SIZE_KGROUP_SIZE_M)
num_stages	num_warps)tritonConfig).0BMBNBKsws         /home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/flashinfer/triton/kernels/sm_constraint_gemm.py
<listcomp>z&matmul_get_configs.<locals>.<listcomp>   s                 	 " " " !	  		
 		
 		
          r    r   r   r   matmul_get_configsr      s%         r   c                    i }|d         |d         |d         }}}|j          d| d| d| d|d<   d	|v r|d	                                         }n|d
         rdnd}d|z  |z  |z  |d|dz   <   |||z  ||z  z   ||z  z   z  |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_elems           r   _matmul_launch_metadatar0      s    
C3icDI!qA[88a88Q88A888CK$g3355"<07a(+a!aC$"$$%!QUQU]QU%:;CLJr   c                 d    | |z  }||z  }t          ||z
  |          }|| |z  z   }| |z  |z  }	||	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_ns
             r   _compute_pidr<   &   sT    **H\)Ky;.==L7\12E''L8E%<r   )r   r    r!   )configskey)launch_metadatar   r   r   r   r6   c           
         t          j        d          }t          j        ||          }t          j        ||          }t          j        ||          }||z  }||z
  }t          j        d|          }||z  }t          j        |||d          D ]}t          |||||          \  }}||z  }||z  }|t          j        d|          z   } |t          j        d|          z   }!t          j        | |k     | d          } t          j        |!|k     |!d          }!t          j        t          j        | |          |          } t          j        t          j        |!|          |          }!t          j	        ||ft           j
                  }"t	          |          D ]}#|#|z  t          j        d|          z   }$| | d d d f         |z  |$d d d f         |z  z   z   }%||$d d d f         |z  |!d d d f         |	z  z   z   }&t          j        |%|d d d f         ||#|z  z
  k     d          }'t          j        |&|d d d f         ||#|z  z
  k     d          }(t          j        |'|(|"          }"||z  }t          |||||          \  }}||z  t          j        d|          z   })||z  t          j        d|          z   }*||
|)d d d f         z  z   ||*d d d f         z  z   }+|)d d d f         |k     |*d d d f         |k     z  },|"                    |j        j                  }-t          j        |-||t          j        |+|,          z            }-t          j        |+|-|,           d S )	Nr   axisTflattendtype        maskotherrI   )tl
program_idcdivarangeranger<   wheremax_contiguousmultiple_ofzerosfloat32loaddottorF   
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   r6   	start_pidr5   	num_pid_nk_tiles	num_tiles	tile_id_coffs_k_for_maskr4   r3   r:   r;   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cs.                                                 r   gemm_kernel_persistentr|   0   s   4 1%%%I<((I<((Iga&&GI%I G#Ii<00O#i/8Iy'4HHH *) *)#%y,
 
u ,&,&BIa666BIa666(7Q;33(7Q;33#BN7L$I$I<XX#BN7L$I$I<XXhl;2:NNN.. 	4 	4B,&1l)C)CCF4 9,vdAAAg/JJF qqq$w)+gdAAAg.>.JJF _T111W5B<M8MMUX  A _QQQW5B<M8MMUX  A &A{33KKW	#'L'
 
u ,&1l)C)CC,&1l)C)CCWQQQW%555	GDRSRSRSGDT8TT!!!T'"Q&747+;a+?@NN5;122F1eTBGF$@$@$@@AA
(((((U*) *)r   EPILOGUE_SUBTILEc                 @   |j         j        }t          j        d          }t          j        ||          }t          j        ||	          }t          j        ||
          }||z  }t          j        | ||g|dg||
g          }t          j        |||g|dg|	|
g          }t          j        |||g|dg||s|	n|	dz  g          }||z
  }||z  }t          j        |||d          D ]3}t          |||||          \  }}||z  }||	z  }t          j        ||	ft          j	                  }t          |          D ]P}||
z  } |
                    || g          }!|
                    || g          }"t          j        |!|"j        |          }Q||z  }t          |||||          \  }}||z  }#||	z  }$|rt          j        ||d|	dz  f          }%t          j        |%d	          }%t          j        |%          \  }&}'t          j        |&|||
                    |#|$g          z            }&t          j        |'|||
                    |#|$|	dz  z   g          z            }'|&                    |          }(|                    |#|$g|(           |'                    |          })|                    |#|$|	dz  z   g|)           t          j        ||||
                    |#|$g          z            }|                    |          }*|                    |#|$g|*           5d S )
Nr   rA   r&   )shapestridesblock_shaper'   TrC   rE   )r   r'   r&   )rF   rY   rL   rM   rN   make_tensor_descriptorrP   r<   rT   rU   rV   rW   TreshapepermutesplitrZ   rX   r[   )+r\   r]   r$   r   r    r!   rd   re   r   r   r   r   r}   r6   rF   rf   r5   rg   rh   ri   a_descb_descc_descrj   r4   r3   r:   r;   rn   ro   rp   rq   rr   ru   rv   rw   rx   accacc0acc1c0c1r{   s+                                              r   !gemm_kernel_descriptor_persistentr      s   " K"E1%%%I<((I<((Iga&&GI%I&!fA!<0	  F &!fA!<0	  F &!fA 0GLLla6G
	  F G#I#i/8Iy'4HHH &0 &0#%y,
 
u ,&,&hl;2:NNN.. 	6 	6B,&FWf-..AWf-..A&AC55KKW	#'L'
 
u ,&,& 	0*[<LA<M*NOOC*S),,C#JD$6$tfkk7G:L.M.M'MNND6eTFKK'LTUDU:U0V$W$WW D BLL'7+R000BLL'7\Q->#>?DDDD&UD6;;7I+J+J$J K u%%ALL'7+Q////M&0 &0r   c           	      V   t          j        d          }t          j        ||          }t          j        ||          }||z  }||z  }||z  }t          ||z
  |          }|||z  z   }||z  |z  }||z  }||z  }|t          j        d|          z   }|t          j        d|          z   }t          j        ||k     |d          }t          j        ||k     |d          }t          j        t          j        ||          |          }t          j        t          j        ||          |          }t          j        d|          }| |d d d f         |z  |d d d f         |z  z   z   } ||d d d f         |z  |d d d f         |	z  z   z   }!t          j        ||ft           j	                  }"t          dt          j        ||                    D ]~}#t          j        | |d d d f         ||#|z  z
  k     d          }$t          j        |!|d d d f         ||#|z  z
  k     d          }%t          j        |$|%|"          }"| ||z  z  } |!||z  z  }!|"                    |j        j                  }&||z  t          j        d|          z   }'||z  t          j        d|          z   }(||
|'d d d f         z  z   ||(d d d f         z  z   })|'d d d f         |k     |(d d d f         |k     z  }*t          j        |&||t          j        |)|*          z            }&t          j        |)|&|*           d S )Nr   rA   rE   rG   rH   rK   )rL   rM   rN   r2   rO   rQ   rR   rS   rT   rU   rP   rV   rW   rX   rF   rY   rZ   r[   )+r\   r]   r$   r   r    r!   r^   r_   r`   ra   rb   rc   rd   re   r   r   r   r   pidr5   rg   r4   r7   r8   r9   r:   r;   rl   rm   rn   ro   rr   rs   rt   rp   kru   rv   r{   rw   rx   ry   rz   s+                                              r   gemm_kernelr      sp   2 -Q


C<((I<((I#i/&&H\)Ky;.==L3-.E##4El"Gl"G	!\222G	!\222Ghw{GQ//Ghw{GQ//Gw E E|TTGw E E|TTGYq,''Fgaaag&2VD!!!G_y5PPQFfQQQWo	1GD!!!G4Dy4PPQF(L,7rzJJJK1bga..// + +GFaaa1q<7G3G!GsSSSGF41q<7G3G!GsSSSfQ;//,**,**u{-..Al"RYq,%?%??Gl"RYq,%?%??GYD!111Iaaa@P4PPFaaag"wtQQQw'7!';<F
q%V < < <<==AHVQV$$$$$$r   )r   triton.languagelanguagerL   r   r0   jitr<   autotune	constexprr|   r   r   r   r   r   <module>r      s+           (
 
 
         3444L) ,L)  ,!L)" ,#L)$ ,%L)& \'L) L) L) 54	 
L)^ 3444X0 ,X0 ,X0 ,X0 ,X0 lX0 \X0 X0 X0 54X0x      3444<% ,<%  ,!<%" ,#<%$ ,%<% <% <% 54	 
<% <% <%r   