o it*@stddlmZddlZddlmZGdddeZGdddeZGdddeZGd d d eZGd d d eZ Gd ddeZ GdddeZ GdddeZ de fddZde fddZejejejde jfde de dededededed ed!e de fd"d#Zd$ejjjjfd%d&ZGd'd(d(eZd)ejdefd*d+Zd,ejd)ejd-ede fd.d/Zd0ejdej fd1d2Z!dS)3)IntEnumNc@eZdZdZdZdS)MajorrN)__name__ __module__ __qualname__KMNr r T/home/ubuntu/vllm_env/lib/python3.10/site-packages/flash_attn/cute/mma_sm100_desc.pyrrc@r)ScaleInrrN)rrrOneNegr r r r rr rc@r)SaturaterrN)rrrFalse_True_r r r r rr rc@eZdZdZdZdZdS)CFormatrrN)rrrF16F32S32r r r r rrc@r) F16F32FormatrrrN)rrrrBF16TF32r r r r r%rrc@r)S8FormatrrN)rrrUINT8INT8r r r r r+r rc@ eZdZdZdZdZdZdZdS)MXF8F6F4FormatrrN)rrrE4M3E5M2E2M3E3M2E2M1r r r r r"0 r"c@seZdZdZdZdZdZdS)MaxShiftrrrr#N)rrrNoShift MaxShift8 MaxShift16 MaxShift32r r r r r,8s r,returncCs~|tjurtjS|tjurtjS|tjurtjS|tj ur tj S|tj ur(tj S|tj ur0tjS|tjur8tjStd|)zJ Map a CUTLASS scalar class to the 3-bit encoding for Matrix A/B. z)Unsupported CUTLASS scalar type for A/B: )cutlassInt8rr Uint8rFloat16rrBFloat16rTFloat32r FloatE4M3FNr"r& FloatE5M2r' TypeError cutlass_typer r r to_UMMA_formatDs       r=cCs>|tjurtjS|tjurtjS|tjurtjStd|)zG Map a CUTLASS scalar class to the 2-bit accumulator encoding. z1Unsupported CUTLASS scalar type for accumulator: ) r2r5rrFloat32rInt32rr:r;r r r to_C_format]s   r@FMNa_majorb_majora_negb_negc_sat is_sparse max_shiftc CsTtt|} tt|} tt|}|dvrtd|dks&|dks&|d@r*td|d?}|d?}d }|d O}|t| d @d >O}|t| d @d>O}||d@d>O}|| d@d>O}|| d@d >O}|t|d @d >O}|t|d @d>O}|t|d @d>O}|t|d @d>O}||d@d>O}||d@d>O}|t| d@d>O}|d@S)u Build the 32-bit instruction descriptor for Blackwell MMA. All matrix/accumulator **types must be CUTLASS scalar classes** – passing integers is forbidden. )@zM must be 64, 128 or 256rLu.N must be a multiple of 8 in the range 8…256r$r#rrr ?l)intr=r@ ValueError)a_typeb_typec_typerArBrCrDrErFrGrHrIa_fmtb_fmtc_fmtm_dimn_dimdescr r r make_instr_descos0   rdopc Cs`t|j|j|j|jd|jd|jtjjj j j krt j nt j |jtjjj j j kr,t j St j S)Nrr)rda_dtypeb_dtype acc_dtype shape_mnk a_major_modecutenvgputcgen05mmaOperandMajorModer rr b_major_mode)rer r r mma_op_to_idescsrqc@r!) LayoutTyperrrr$N)rrr SWIZZLE_NONESWIZZLE_128B_BASE32B SWIZZLE_128B SWIZZLE_64B SWIZZLE_32Br r r r rrr+rrswizzlecCst|}||dd|d}dd|dD\}}}|dkr8|dkr+td tjtjtjtjd |S|d krI||fd krFtd tj Std)NcSsg|]}t|qSr )rY).0xr r r sz _layout_type..,r$r#u/Unexpected swizzle shift – want S==3 for M==4)rrrr#r%)rrz.Only Swizzle<2,5,2> supported for 128B_BASE32Bz3Unsupported swizzle triple for UMMA smem descriptor) strindexsplitrZrrrtrxrwrvru)ryswz_strinsideBrASr r r _layout_types$ rlayoutmajorc Cs t|}d}d}d}tjdtjdtjdtjdtjdi|}|tjur|tjur)dnd}t |||f} t | ds=t d| j dd} |tjurQ| dkrQt d| j dd} | |kr`t d| j dd| j dd} } |tjurz| | }}n^| | }}nX|tjkrt d t |jdddkst d t |d } t | dst d | j dd} | |krt d | j dd} |tjur| dkrt d | j dd} | | }}d}||d@d>O}||d@d>O}||d@d>O}||d@d>O}||d@d>O}|t|d@d>O}|d@S)z Convert a 2-D *shared-memory* Cute layout into the Blackwell 64-bit smem-descriptor, without the smem start address. layout must correspond to layout of an uint128 tensor. rrrr$rM)rrrz9Not a canonical UMMA_MN Layout: Expected profile failure.z8Not a canonical UMMA_MN Layout: Expected stride failure.z+SWIZZLE_128B_BASE32B is invalid for Major-Kz>Not a canonical UMMA_K Layout: Expected MN-size multiple of 8.)rMrz8Not a canonical UMMA_K Layout: Expected profile failure.z7Not a canonical UMMA_K Layout: Expected stride failure.i?rS r#.rN14=l)rrrrtrxrwrvrurr rklogical_divide is_congruentrZstridesizeshaperY)rryr layout_typeVERSIONLBO_MODE BASE_OFFSETswizzle_atom_mn_sizeswizzle_atom_k_sizecanonical_layout stride_00 stride_10 stride_01 stride_11stride_byte_offsetleading_byte_offsetrcr r r make_smem_desc_basesd         r start_addrcCs|d@d?S)Nir$)toint)rr r r make_smem_desc_start_addrsr)"enumrr2 cutlass.cuterkrrrrrrr"r,rYr=r@rrr-boolrdrlrmrnMmaOprqrrSwizzlerLayoutrPointerr?rr r r r sX        4 I