o
    i                     @   sv   d dl Z d dlmZ d dlmZ ej		 	ddejdejdejdejde je	 d	e je
 d
e je	 ddfddZdS )    N)	warpgroupF	tiled_mmaacctCrAtCrB	zero_initwg_waitswap_ABreturnc           	   
   C   s   t |rt| |||||dd d S t  t| j}|tj	j
|  t t|jd D ]}t|||d d |f |d d |f | |tj	j
d q0t  t |dkr`t| d S d S )NF)r   r   r	      Tr   )cutlass
const_exprgemmr   fencecutemake_mma_atomopsetField
ACCUMULATErange_constexprsizeshapecommit_group
wait_group)	r   r   r   r   r   r   r	   mma_atomk r   T/home/ubuntu/vllm_env/lib/python3.10/site-packages/flash_attn/cute/hopper_helpers.pyr      s   
&r   )Fr   F)r   cutlass.cuter   cutlass.cute.nvgpur   jitTiledMmaTensor	Constexprboolintr   r   r   r   r   <module>   s0   	