ukiddlmZddlZddlZddlZddlZddlmZmZm Z m Z ddl m Z ddl mZddl mZddl mZddl mZddlZd d lmZd d lmZd d lmZd d lmZdZdZdZej<Zej@Z!ejDZ#ejHZ%ejLZ& d9 d:dZ' d;dZ(ddZ, d? d@dZ-dAdZ.dBdCdZ/dDdZ0dBdEdZ1dFd Z2dGd!Z3dHd"Z4dId#Z5ejld$Gd%d&ejnZ8 dJd'Z9dKd(Z:dLdMd)Z;dLdMd*ZdPd-Z?dPd.Z@ejld$Gd/dZA dQd0ZBdRd1ZCdRd2ZDdSd3ZEdSd4ZFdRd5ZGdRd6ZH dT dUd7ZI dT dUd8ZJy)V) annotationsN)AnyCallableIteratorcast)ir)arith)llvm)memref)nvvm)fragmented_array) mma_utils)utils) LaunchContextlFctjj}tjj}tjj d} d} |d|cxkrdksJJ| |z} | dz} ||k(rd} n||k(rd} n|| k(rd} nt d|| | dzz} ||k(r |||hvsJd} n|tjjk(r ||k(sJd} n|tjjk(r |||hvsJd} nf|tjjk(r |||hvsJd} n:|tjj dk(r || k(sJd} nt d || | d zz} | | d zz} | |d zz} | |d zz} |dzs|dkDrtd|| |dz dzz} |d zs|dkDrtd|| |dz dzz} tjtjj d| S)N rr zUnsupported accumulator dtype: Unsupported input dtype:  +N must be a multiple of 8 and <= 256, got: ,M must be a multiple of 16 and <= 256, got: ) rF16TypegetF32Type IntegerType get_signlessNotImplementedErrorBF16TypeFloat8E4M3FNTypeFloat8E5M2Type ValueErrorr constant) mn acc_dtype input_dtype transpose_a transpose_bsparsity_selectorf16f32i32desc d_type_val ab_type_vals ^/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/experimental/mosaic/gpu/tcgen05.pycreate_instr_descriptorr=/sY # # ##B'# $" ! %A %% %% % DFND#JCJCJ  ? {K LL:?$C c "" "Kbkkoo''   Kb))--// c "" "Kb''++-- c "" "Kbnn11!44   K  9+G HH;! $;" $+ $+ $Ua#g B1#F GG16b.$Vq3w CA3G HH16b.$ 33B7 >>c Ld} d|cxkrdksJJ| |dzz} | ||dzz} | ||dzz} | |dzz} | |dzz} |dzs|dkDrtd || |d z d zz} | tjjk(rd } n2| tjjk(rd} nt d | | | dzz} |dzs|dkDrtd|| |dz dzz} | |dzz} t jtjjd| S)Nrrrrrrrrr rr!r zUnsupported scale type: rr"r) r-rFloat8E8M0FNUTyper%r+r)r r.r'r() get_input_encodingr/r0a_typeb_type a_scale_idx b_scale_idxr3r4 scale_typer9scale_encodings r<_create_scaled_instr_descriptorrKmsh $ k A   + $ V $ ))$ V $ **$+ $+ $Ua#g B1#F GG16b.$2''++--NR((,,..N  8 E FF.B $WC CA3G HH16b.$+ $ 33B7 >>r>c&d}t|g|i|S)Nc|tjjk(ry|tjjk(ryt d|)Nrr r)rr+r%r,r)tys r<rDzAcreate_scaled_f8f6f4_instr_descriptor..get_input_encodingsI R $ $ &&  r  $$& &  ";B4 @ AAr>rKargskwargsrDs r<%create_scaled_f8f6f4_instr_descriptorrTs B )); Md Mf MMr>c&d}t|g|i|S)Ncb|tjjk(rytd|)Nr r)rFloat4E2M1FNTyper%r)rNs r<rDz=create_scaled_f4_instr_descriptor..get_input_encodings/ R $ $ &&  ";B4 @ AAr>rPrQs r<!create_scaled_f4_instr_descriptorrXs B )); Md Mf MMr>T) a_swizzle b_swizzlea_scaleb_scalea_sparse_metadata accumulate collectiveTMEMRefclH|dk(s|dk(r tdtjjd} tjjd} t |t r3t jtjjd|}| rdnd} |dux} |duk7r td|du}| r |r tdtjj |jstd |jtj|\\}}Ht |tr|j\}}|j}| r td |d k7rtd |dt!j"Hz}t%|j| | }|j&|k7rwtd|d|j&tjj |jstd|jtj|\\}}}|r|dz}||k7rtd|d|H|k7rtd|dH|j||| zfk7r td||| zfd|j|d k(r9|j&t)d x}k7rtd|d|j&d}n|dk(r| r td|r tdt%|j| d }|j&|k7rtd|d|j&| rd}n@d}d|zt!j"Hz||zkDrtd|dtd|tj*j-}tj.j-}tjjd}H|k(s!Htj0j-k(rNH|k(r |r td| rtdHd|j|k7rtdHd|jH|k(rL| rtdHd|j|k7r|j|k7rtdHd |jt3Hfd!tj4tj6hDre|j|k7r*|j|k7rtdHd |j| r|j|k7rtd"Hd#|jt3Hfd$tj8hDrH|r td%| stdHd&|j|k7rtd"Hd#|jHtjjdk(r8| rtdHd|j|k7r&td'|jtd(H|}dt;|d|zz|zt!j"Hz}|r|dkrd}d}| r>|jtj<j-k(rdnd}t;|d)|z}| rdnd}| rd*nd+}|j|k(r |dz}|d,z }|dkDr |d-|d.z }||z|zd/k7rtd0|d1|d2||s| r!|j?dk7rtd3||d4kDr!|j?dk7rtd5|| r|d kDr td6|d7kDr td8tA||zd4| z} ||zrtd9|d:|||zrtd;|d:||| zrtd<| d:|||z}!||z}"|| z}#Htj*j-k(rtjBj-nH}$d}%| rE|d k(sJ|dzrtd=|||J|j}%|jtj<j-k7rC|jtj6j-k7rtd>|j|j|jk7r%td?|jd@|j|j|||zfk7r!tdA|dB||zdC|j|j|| z||zfk7r!tdD|dB||zdC|j|rtEt|}|dzrtdE||j||dzfk7r tdF||dzfd|j|jtjjdk7rtdG|jt |ts.tjF||||d|zzfdHddIJ\\}&}'\}(})}*n"tjHjJ}*d}'dx}(x})}&tjF|||| fdI|rdnddIJ\\}+},\}-}.}/| rht!j"|$d)k(rP|*tjHjJk7r tdK|/tjHjJk7r tdL|r-|dk(r(|/tjHjJk(r tdMt jtjjdd}0| | z}1|#|z}2|j&jLd/d)zd/k(sJ|j&jLd/d)z}3|r |jNnd}4| r |jNnd}5| r |jNnd}6tQjR|!|#|"D]\}7}8}9t |trU|!dk7r tdN|d|zz}:|jUtUdt!jV|9|:z|:jN};n+|7|(z|9|)zz}<|&d/|&dtjX|<zf};|8|-z|9|.zz}=|+d/|+dtjX|=zf}>|4R|#dk7s|!dk7r tdO|dzd/k(sJ|dz}?t jZ|4t!j\|9|?z| }@nd}@|5|6|!dk7r tdP|#dk7r tdQ|J||d)zzd/k(sJ|dzd/k(r| dzd/k(sJ||d)zz}At jZ|5t!j\|9|Az|zdz| }Bt jZ|6t!j\|9|Az|1zdz| }Cndx}B}C|9d/k(r|n|0}D|8|2z|8|2z}F}E|E|3zdz|F|1zz}G|!dk7r tdRt_t jZ|jNt j| G|;|>f|j|| || |*tjHjJk7|/tjHjJk7|'|,BC@D|$|%dSy)TNrzNo swizzle is not supportedr@r rz-Either none or both scales should be providedz'Block-scaled sparse matmuls unsupportedzB must be a memref, got: z.A in TMEM unsupported for block-scaled matmulsrz:Only M=128 is supported for MMA with A in TMEM, but got M=packingzA layout mismatch: expected z, got zA must be a memref, got zFMMA requires A and B to have the same contraction dimension (K), got: z and z9MMA requires A and B to have the same element type, got: z%Accumulator shape mismatch: expected z&Accumulator layout mismatch: expected z0MMA with block scaling is not supported for M=64z!Sparse MMA not supported for M=64rzSwizzle=z/ is too big for MMA with M=64. Try lowering it.z5Only M=128 and M=64 are supported for MMA, but got M=zSparse MMA unsupported for f32zMMA with element type z does not support block scalingz2 only supports accumulators of type f32, but got: z9 only supports accumulators of type f32 or f16, but got: c3@K|]}|jywN isinstance.0t element_types r< zmma..(s"  ll<  z#Block-scaled MMA with element type z* only supports f32 accumulators, but got: c3@K|]}|jywrfrgris r<rmzmma..6s %&all<  rnz#Sparse MMA unsupported for f4e2m1fnz only supports block scalingzBMMA with element type s8 only supports s32 accumulators, but got: zUnsupported element type: rz2 CTAz1 CTAz integerz with z lane groupsrzIn z MMA, N must be a multiple of z, got N=zOOnly N that is power of 2 supported for sparse and block-scaled MMA, but got N=rz.The only supported N > 256, is 512, but got N=z.Only N <= 128 are supported for collective MMArz#Only N <= 512 are supported for MMAzM must be a multiple of z, got: zK must be a multiple of zN must be a multiple of z>MMA with block scaling requires N to be divisible by 32, got: z KJA'$&&iG4,?@ D EEt+)9 G HH  ! !!&& ) 09 ::"55a8&1a,7 GGEArGGM  :  Cx "\]^\_ ` aaU^^L99( %5O xx?"  ((9z J  == # #AFF + 1!&&: ;;&99!<GQ]!GB"W  E!  ]"  ?% ~ / WWAK   /AK0@/AyQ #Xxx)%%&WWI /   +-+>+>*?   E FF   "<.0L M  ww#~  / ~>%%&WWI / r~~22155  "<.0O P  ww#~  ggY    :<.I JJ-c)q9}5yAAU^^T`Eaa-=2%M + )=)=)A)A)CC"K q;7M&bA#g)WW^ IQ 6- 55I=--2  i[67H6IJ  9!++-1"4  S  W!#  NqcR SSAG E FF3w : ;;a=(#.9- / gaSI JJ / gaSI JJ / gaSI JJ - ( - ( - (!- 0@ @bl  8O82v  J1# N   7#66 6  --1133 MMR00446 6  H X }} %  -gmm_=mm_  }}A,--  .qcA4D3EFmm_  }}Wa;&677  .qcA4D3EFmm_  W&782v KA3O PP1a1f+-  7AF }E%% & ( ".."="=a"@@  %% & (  Aw  ## !=Q]#CD   (', + I8<<<'+ !!/%b2  &{%*)5>>"23q8IMMOO#  W IMMOO#  W B9 7  H  33A6 :$*W4]*,  ! !! $q (A -- -hh..q1Q64=(004)2goo)2gooJJx8<Bjb"b!W Q!"NOO%!i-8o WWU4[%((2+?"Q R Z Zd&&.>)>>h!nk!ny/D/DX/NN Od$$r,<'<c tjjd}tjjd}tjjd}|xsd\}}|\}}tdt j |xsd|DsJ|du|duk(sJ|du}| du}t j| }d|zdzdz|z}d|z}d}|r0|rJtjj| stjj| rO|tjjk7rtd|d }d }tjt | }ntj"j| rm|s|rJtjt$| }|tjjk(rd }d }n4|tjjk(rd}d}nt'd| d}d} ntj(j| stj*j| rd}ntjj| rd}natjj| rd}n?tjjdj| rd}nt'd| dx} }d}|rd nd}!|du}"|"rdnd}#|rdnd}$|rdnd}%| |rdndz } d}&d}'d7d}(t-| |zD]E})|r|J|rJd |z}*|)|z|*z}+|| |!z| |!z| | |+|+||},| dk(sJ| dzdk(sJt/j0||)|zd z}-t/j0||)|z| zdz|!z}.t/j2||-t/j2||.f}'nvd}/|rZd|cxkrdksJJ|}0d|0z}1| |z|1zdk(sJ|)|1z}/t/j2| t j4|)|1zd z|f}&t7| |!z| |!z|| |||/ },|"r]||zd|zz}2|)|2z}3t|tj8sJ|j:tjjdk(sJ|}4n||J|\}4}3|3|(|)||z }3|\}5}6|6|(|)||z }6|3d!z|3dz }8}7|6d!z|6dz }:}9t=j>tj@jCd"||4|5|,|g|'|&d#|"rd$nd%d&|7d'|9d(|:d)|"rd*nd+|8d,d-|$d.|!d/d0|#d1|%d2|d3d4| zd56t/j0|d}Hy)8Nr rrbNNc3,K|] }|dzdk(yw)rrN)rjss r<rmz_do_mma..Ks RQQVq[ Rrrz5Scale element type mismatch: expected f8e8m0fnu, got z"mxf8f6f4.block_scale.scale_vec::1Xr)rIzmxf4.block_scale.scale_vec::2Xrz"mxf4nvf4.block_scale.scale_vec::4Xz,Unsupported element type for block scaling: z [$5], [$6], z,r,rr6f8f6f4i8z Unsupported input element type: ctrf)r))rRs r<create_scaled_instr_descriptorz/_do_mma..create_scaled_instr_descriptor~s r>z[a_desc]a_descz.spz[$5], z,rct|dzt|k(sJg}|D]}|j||z||z}|j|tdt||dD}|dz S)Nr c3,K|] \}}||zywrfr)rjirs r<rmz/_do_mma.._get_offset..sC41aQCrT)strictr)lenappendsumzip)idx idx_tilingstridesidxsrkoffsets r< _get_offsetz_do_mma.._get_offsets{ z?Q #g, .. . D  kk#( !Gc KK C3tWT#BC CF Q;r>rr)r5l !llvm.voidz[{ .reg .b32 a_desc_low, a_desc_high, b_desc_low, b_desc_high; .reg z.b32z.b64zK a_desc; .reg .b64 b_desc; add.s32 a_desc_low, $1, z&; add.s32 b_desc_low, $2, z+; mov.b64 b_desc, {b_desc_low, z}; zmov.b32 a_desc, a_desc_low;zmov.b64 a_desc, {a_desc_low, z};z tcgen05.mmaz .cta_group::z.kind::z [$0], z , b_desc, z$3, z $4; }z r,r,r,r,bThas_side_effects)rintrtuple[int, ...]rr)"rr'r(all itertoolschainrrr,rhr+rCr%r- functoolspartialrTrWrXr)r$r*ranger r.rrr=Valuerr inline_asmTypeparse);d_addra_desc_or_addrb_descrwrxryrzr{r|r}r/r0rvrlr~rur^r_i1r8ra_k_idx_tilingb_k_idx_tilingrr elem_bitwidthinstr_krd scale_stepskindr extra_ptxextra_constraintsr a_in_tmema_ptx sparse_modsparse_meta_ptx sparse_addr scales_addrsrk_stepscale_vec_widthscale_idi_desca_scale_addr_offsetb_scale_addr_offset sp_selectorselector_widthk_steps_for_col_incrra_enc_addr_baseb_enc_addr_baser a_offset_low a_offset_high b_offset_low b_offset_highs; r<rr2s1( ~~""1%" ##B'# ##B'# + ;|.+ +.+ R)//+2C["Q R RR R $ LD$8 99 9$&)4')...- ]a " $ 5' ] "'+= $$\2 &&|4 r33779 9CDVCW X  2dk'0'8'8 /**h  11 1  BNN$?$?$C CC C&o  'K,CC C"0ox+fnkBBh &OX FNK@@H"*Z"7R-L"*Z"7R-LOO  l# /6:c cWbc %&623%%1N3%%1N3++8/: .7 *?]^k]llo=p qr"|<yvWUGS]^m]nnrs|r}~   '' A&JKE'r>c*t|tjr|j}nD|jt j jdk7rtd|j|r| td|jdk7r tdt jjd}tj|d}tj |tj"j$|ytj |y) N !llvm.ptr<3>z9barrier must be a Mosaic barrier or a SMEM pointer, got: z,ctx must be provided for collective barriers)rr r z:Collective arrivals only support (2, 1, 1)-shaped clustersrr)groupmulticast_mask)rhr BarrierRefget_ptrrrrrr- cluster_sizer)r'r(r r.r tcgen05_commit CTAGroupKindCTA_2)barrierr_ctxi16masks r< commit_arriver+s ))*ooG||rww}}^44  LL>   { E FF 9$  \ ]] .. % %b )C >>#q !Dt((..t  r>c|r0|jdk7s d|cxkrdksntd||Stdd|dz jz}|dkDrtd|d|S)a!Returns the exact number of columns to allocate in TMEM. The number of columns is rounded up to the nearest power of 2. Args: ncols: The number of columns to allocate. exact: If true, throws an error if the number of columns is not a power of 2 and within [32, 512]. r rrz6ncols must be a power of 2 and within [32, 512], got: zAfter rounding up, got z$ columns, exceeding the limit of 512)rr-r bit_length)ncolsexacts r<tmem_alloc_exact_ncolsr0s  AR5%7C%7 OPUwW XX , A%!)//11 2E s{  #E7*N O  ,r>ctjj|jrtj|j}|jtj j dk7rtd|tj|std|tj|jdk7rtd|tj|d}nD|jtjjdk7rtd |jt!||}|rt"j$j&nt"j$j(}tj j d}t#j*|tj,||| |fS) Nrz&tmem_addr must be an i32 memref, got: z)tmem_addr must be in shared memory, got: r z.tmem_addr must contain a single element, got: r) memory_spacerz4tmem_addr must be an SMEM pointer or a memref, got: r)rrrhrrlr'r(r-r is_smem_refmathprodr memref_ptrrrr0r r%r&CTA_1 tcgen05_allocr) tmem_addrr.r_r/ref_tyrr8s r< tmem_allocr< sL]]inn- ]]9>> *F bnn99"== ?xH II   V $ B6(K LL yy!# GxP QQ  ;I~~~66 KINNK[\ ]]  .%%/$   ! !T5F5F5L5L% ##B'#  Iuwwuc':% H% OOr>c|jtjjdk(sJtjj d}t j||S)Nrz !llvm.ptr<6>)rrr'r(rrr inttoptr)r:ptr_tys r<_tmem_addr_to_ptrr@sF 2>>66r: :: : 77== (& vy ))r>c|jtjjdk7rt d|jt ||}|rt jjnt jj}tjjd}t jt|tj|||y)Nrztmem_addr must be an i32, got: r3)rrr'r(r-r0r r%r&r8tcgen05_deallocr@rr)r:r.r_r/rr8s r< tmem_deallocrC#s^^r~~22266 6y~~6FG HH  .%%/$   ! !T5F5F5L5L% ##B'# "EGGE3$7ur>c|rtjjntjj}tj|y)Nr3)r r%r&r8tcgen05_relinquish_alloc_permit)r_rs r<tmem_relinquish_alloc_permitrF.s1%/$   ! !T5F5F5L5L%&&U3r>c :|jdk7s|dkDrtd||xdk(rd}n xdk(rd}ndk(rd}n td |d ||z}|d kDrtd |d |d|ddjdt |D}d|zdz}||fS)Nr rz*num must be a power of 2 and <= 128, got: 32x32b16x128br16x256brzshape=z is unsupportedz!TMEM translation too big : shape=z and num=z involve z5 registers per-thread, which exceeds the limit of 255,c3&K|] }d| yw)$Nr)rjrs r<rmz&_tmem_access_helper..Ds:Q1QC:s{})rr-r)joinr)rnumnum_regs regs_vectors r<_tmem_access_helperrU2s]]_S3Y A#G HH h h h 65(/ : ;; c/( ^  ,eXZ3&9 :J L :%/::+k!C'+ ; r>ctjjd}t||\}}|rdnd}t j tj jddjdt|Dzdz|gd|d ||d |d |d d |zdzd}t|D cgc]} t j||| gc} Scc} w)Nrz .pack::16brz!llvm.struct<(rLc3 K|]}dyw)r8Nr)rj_s r<rmz_tmem_load..Os%Ie%Is z)>ztcgen05.ld.sync.aligned..xz.b32 z, [$z];z=r,rTr) rr'r(rUr rrrrQr extractvalue) r:rrRpackr8 num_out_regsrTpack_modregsrs r< _tmem_loadr`Is ##B'#1%=, !\r( ggmm SXX%IU<5H%II ID Pk r#xjk]$|n\^_ lS  $6;<5H I$  Cs + II Is4Cct||\}}|rdnd}tjtjj dg||d|d||d|d|d d |zd zd y) Nz .unpack::16brrztcgen05.st.sync.aligned.rYz.b32 [$z], ;zr,rZTr)rUr rrrr)r:rrRr_unpackr]rTr^s r< _tmem_storerdYsz1%=, %^2(//ggmmL!i r#xj ~SQ\P]]^_ \C r>)frozenc<eZdZdZddZddZdfd Zd dZxZS) TMEMLayoutaRepresents the way a shape is laid out in TMEM. The layout describes how the shape is split across the 128 rows (lanes) of TMEM. We reinterpret warp_dims as the partitioning of TMEM into 4 banks, each accessible from a single warp. The 32 lanes inside each bank are assigned consecutive elements from lane_dims. The data within each lane is linearized in row-major order, with each vector padded up to 32 bits (wider vectors are unsupported). c"t|dk7rtd|tdt||jDrt|d|j|j dd|zx}hvrtd|d|d |j y) Nr'TMEM can only represent 2D shapes, got c3,K|] \}}||zywrfr)rjrrks r<rmz(TMEMLayout.check_type..ts >TQ1q5 >rz& is not divisible into tiles of shape r rzFor z+-bit types, the vector length must be 1 or z , but got: )rr-rrr vector_length)selfrr fully_packeds r< check_typezTMEMLayout.check_typeqs 5zQ @H II >S(<(<= >>  G9$:N:N9O P  !R8^%C\!DD   ^<(:(:'; = Er>cn|j||d}|jD],}t|tjs||j z}.|j D],}t|tjs||j z}.tj|tz|jz|zS)Nr ) rn warp_dimsrhfa Replicatedtimes lane_dimsr5r6 TMEM_ROWSrk)rlrrreplication_factordims r< cols_in_shapezTMEMLayout.cols_in_shape~sOOE8$~~( C 'cii'(~~( C 'cii'( 99U y (D,>,> >AS SSr>ct|}t|j|j|j |j dS)NF)_check_canonical)super canonicalizergtilingrprt vector_dim)rlr __class__s r<r|zTMEMLayout.canonicalizesB W ! #F    r>ctj|j|j|j|j Srf)rq TiledLayoutr}rprtr~rls r<as_tiled_layoutzTMEMLayout.as_tiled_layouts, >> T^^T^^T__ r>)rrrrreturnNone)rtuple[int, int]rrrrrrg)rfa.TiledLayout) __name__ __module__ __qualname____doc__rnrxr|r __classcell__)rs@r<rgrges  Tr>rgc|t|k(rtS|t||k(rtjS|t ||k(r t |Std|d)Nrc TMEM layout  is not supported)rLAYOUTtmem_half_lane_layoutrq WGMMA_LAYOUTtmem_m64_collective_layoutfa_m64_collective_layoutr-) tmem_layoutcolumnsrds r<!_infer_tmem_load_registers_layoutrsa'88 M)'7CC ??.wHH #G ,,\+.?@AAr>c t|dk7rtd||dkDs|jdk7rtd||d|zrtd||dtk(r t |S|dtdzk(r |rt |d|St |d|Std|d td tdzd |dd ) Nrrirr ,Packing must be <= 8 and a power of 2, got: zugF GG q[G%%'1, CG9M NN 1X STYSZ[ \\ 1X w '' Qx9>! 'a' :: "58W 55  eW% ;d9>*/%(1 F r>c|jdk7rtd|ttjt |ftj |ffdddS)zCA TMEM layout used for 1CTA MMA with M=128 and 2CTA MMA with M=256.r z#Packing must be a power of 2, got: )rprtr~rr-rgrqTilingru WARP_SIZErcs r<rrs[ A :7)D EE ii)W% g'>?@  r>c||kDs|jdk7rtd||dzrtd|ttjt dz|ftj dz|dzf|ffddd S) z*A TMEM layout used for 1CTA MMA with M=64.r rr'Columns must be a multiple of 16, got: r))rrrrrrds r<rrs w'++-2 CG9M NN r\ >wiH II ii >7 # <<1 gl + *   r>c|dkDs|jdk7rtd||dzrtd|ttjt dz|ftj |dzf|ffddd S) z+A TMEM layout used for 2CTA MMA with M=128.rr rrrr)rr)rrrrrs r<rrs q[G%%'1, CG9M NN r\ >wiH II ii >7 # <<A & *   r>c|dzrtd|tjtjtdz|ftj |dzfddfddd S) zEThe register layout for transfers to/from tmem_m64_collective_layout.rrr)rrr))rrrr)r-rqrrrur)rs r<rrsh r\ >wiH II ii >7 #bllGqL%A64   r>cttjtdftdzdfftjdfddS)zA TMEM layout for A and B scales in .scale_vec::1X configuration. See https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-mma-scale-factor-a-layout-1x rr )rsrrr)rgrqrrurrrr>r< scales_layoutrsD ii)Q)q.!!456Q')  r>clttjtdftdzdfddfddd S) zA TMEM layout for A sparsity metadata. See https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-sparse-matrices-sparsity-selector-kind-tf32-m128-256 rrr )rr )rr )i)rrrrr)rgrqrrurr>r<sparse_meta_layoutrs< ii)R9>1"5wGH  r>ceZdZUded<ded<ded<ded<edd Zd Ze d dd Zdd Z dddZ ddZ ddZ y )r`ir.Valuerrrir.Typerrgrc.|jjSrf)rrkrs r<rdzTMEMRef.packings ;; $ $$r>ctj|j|jz}|dkst d|y)NrzFExpected packed packed bitwidth to be <= 32, but got: packed_bitwidth=)rrrrdr-)rlpacked_bitwidths r< __post_init__zTMEMRef.__post_init__!sGnnTZZ04<<?O b +)+- .. !r>Nc4tjjd}tjj |j st d|j tj|j }tj|st d||j|k7rt d|tj|jdk7rt d|tjjd}t!j"||g|j$z} |ddkrt d|d|| t d t'||d }n%|j)|tj*||| |||S) Nrz2tmem_addr_ref must be a memref or a pointer, got: z-tmem_addr_ref must be in shared memory, got: z*tmem_addr_ref must be an i32 memref, got: r z2tmem_addr_ref must contain a single element, got: rz+TMEM refs must have at least 32 rows, got: zAcollective argument must be provided when TMEM layout is inferredrc)rr'r(rrhrr-rr4rlr5r6rr ConstantOp create_indexr loadrankrrnr) cls tmem_addr_refrrr_rr8 addr_ref_tyi0r:s r< from_alloczTMEMRef.from_alloc'sv .. % %b )C == # #M$6$6 7 KML^L^K_` aa-- 2 23K   [ ) F{mT UU3& CK=Q RR yy""#q( KK=Y ZZ    & &q )B MB4+2B2B+BCI Qx"} DU1XJO PP ~   O  "%Q?f u~~e45 y% //r>c 0tjjd}tj||j \}}}t |r td|dgt|zk(r|t|j k(r|S|jt|jk7rtd|j|ddk7s |dtk7r td|ddzrtd |dd |d}t|tj st#j$||}|j&tj(j+k(rt#j,||}|jdk7r4t#j.|t#j$||j}t1t#j2|j4|t7t8t:t:ft9||j|j< S) Nrz$TMEM can only be sliced, not indexedrrcz=Slicing only implemented for refs with standard layout, got: z TMEM cannot be sliced along rowsr rz6TMEM column slice length must be a multiple of 8. Got r)rrrr)rr'r(r parse_indicesrrr-rlistrrrdr)rurhrr r.r IndexTyper% index_castdivuir`rrrtuplerr)rlrr8base_idx slice_shape is_squeezedcol_idxs r<rz TMEMRef.sliceKs .. % %b )C).)<)>A3X&&;$tzz:J+J k {{)$,,??  kk]  {a;q>Y6  B CC1~  Q  # qkG grxx (sG,g||r||''))  g.g ||q GU^^C%FGg  4<<15c?E+$67{{jj  r>cv|j}|$t|j|jd|}t j |j }|jt|k(}|j|j}|ddk7r td|tk(rc|jt|k(rJt|j|jd|j |jj|}n||jjk(r||zdk(r t!|j"dk(sJt%d|j&DrJt%d|j(DrJt+j,d |j/dj&D}t+j,d |j/dj(D} || z} t1|j|jd| z|j ||j|}n|t2k(ro|rm|d k(r|dk(s|dk(r^|j4dk(rOt1|j|jd|j |t2j4j|}nv|t6j8k(r|jt;|jd|k(rt|j|jddz|j |} | jdd k(sJt=j>| dd| ddgd }|jj|}n|tA|jdk(r|jtC|jd|k(r]|j|j}t|j|jddz|j |j|}ntEd|jd|dt7jF|||S)Nr rcrzLoading multiple row tilesrrc3PK|]}t|tj ywrfrhrqrrrjrs r<rmzTMEMRef.load..LaZ2==1L$&c3PK|]}t|tj ywrfrrs r<rmzTMEMRef.load..rrc3lK|],}t|tjr |jnd.ywr Nrhrqrrrsrs r<rmzTMEMRef.load...$ 2==1!''q 8$24c3lK|],}t|tjr |jnd.ywrrrs r<rmzTMEMRef.load..rrrraxiszLoads from TMEM layout z to register layout z are not supported) _registers_layout _is_signed)$rdrrrrrrrregisters_shaper)r _load_32xcolsrTreshaperrrrrprtr5r6remove_dimension_load_32xcols_nativeTMEM_NATIVE_LAYOUTrkrqrrr concatenaterrr-FragmentedArray) rlr is_signedrdrhas_default_layout regs_shape registerswarp_split_factorlane_split_factor split_factor raw_registerss r<rz TMEMRef.loadlsllG ~0 ++tzz!}gf~~djj)H(;G(LL'' 3J!}  < == DKK+>w+OO ,, 1 tzz7 ''*  4;;..0 0Wx5G25M '' (A -- -L6;K;KLLL LL6;K;KLLL L))$**1-77$))$**1-77$')::l& ,, 1 5tzz7G  % %*< RGqL Nv33q8& ,, 1 tzz7c.t|tjstd||j|jk7r%t d|jd|j|j |jk7r%t d|j d|jt|jtjstd|jd|j}|jt|k(}tj|j}|jtk(r=|r;t|j |j"j$j'd |y|j|jj)k(r9||zd k(r1t+|j |j"j'd |y|jt,k(r[|rY|d k(r|d k(s|d k(rJ|jj.dk(r1t+|j |j"j'd |y|jtj0k(r|jt3|jd |k(rk|j"j$j'dd }t5j6t5j8|dd d}t|j ||y|jt;|jd k(rY|jt=|jd |k(r2t|j |j"j'dd |yt d|jd|jd)Nz+TMEM stores expect a FragmentedArray, got: zStored array has shape z, but TMEM has shape zStored array has dtype z, but TMEM has dtype zStored array has layout z&, but TMEM stores expect a TiledLayoutrc)rrrrrr rrrrzStoring from register layout z to TMEM layout r)rhrqr TypeErrorrr- mlir_dtyperrrrdrrrr_store_32xcolsrrrrr_store_32xcols_nativerrkrrrrsplitrr)rlvaluerdrrrs r<storez TMEMRef.stores eR// 0 CE7K LL {{djj  #EKK=1jj\   4::%  #E$4$4#56jj\   ellBNN 3 0>de ffllG(;G(LL~~djj)H ||v"4 ,,))11':G 446 67X;MQS;SDLL%//*A*A"*EwO + +0B RGqL Nu||99Q>DLL%//*A*A"*EwO ' KK0AP P//##++Ar2i..)QQ!?aHiT\\9g6 1 1   ++3 1 w T\\5??#:#:1b#A7K  )%,,8kk]+ - r>c xtjjd}|jj |j t j|j}tjt jtj|t j}t|D]}ttj |j"tj||}t%j&|t$j(j*|}t j|j}d|z}|j,dk(rY|dkr3tj.tjj||}t j0||j}ng|j,|k(r@t j0|tj2j5|f|j}nt7d|j,t j8d|d||dy)Nrr Unsupported packing: z[{}, z]: {}F)uniform)rr'r(rrxrrrrr remui thread_idxr.WARPGROUP_SIZErr@rrr tcgen05_ldTcgen05LdStShape SHAPE_32X32Brdtruncibitcast VectorTyper%r) debug_print) rlr8num_colslanerptrvaldtype_bitwidth full_packings r< _debug_printzTMEMRef._debug_prints .. % %b )C{{((U^^DJJ5OPH ;;u'')5>>#u?S?S+T UD 8_H ejju~~c17MN Oc OOC!6!6!C!CS Ic~~djj1n>)l   B  R^^88H#N#mmC, <<< 'mmC!2!2L?DJJ!OP!$9$,,"HII '!G,dCGHr>)rrr) rrrrr_ bool | NonerzTMEMLayout | Nonerr`)rr`)rzfa.TiledLayout | Nonerrrfa.FragmentedArray)rrrr) rrr__annotations__propertyrdr classmethodrrrrrrr>r<r`r`s   .  % %.  !%"& !0!0!0  !0  !0!0!0FB@D2hHr>c#Ktjjd}|\}}||zdk(sJ||z}|d|dztj|zz} d|zdk(sJd|z} d|j dz z} t | d| z} t| D]} tj|tj| |zdz|} d}| }||z x}dkDrg||kDr |dz}||kDr t|||z}tj| tj||z|z|}||| |f||z }||z x}dkDrg||k(rJyw)a$Generates a sequence of parameters for a given TMEM read or write. Arguments: base_addr: The base address of the TMEM region. cols: The number of logical columns to transfer. atom_shape: The logical shape of the tile written by the warp in a single TMEM transfer. tmem_packing: Packing degree in TMEM. When packing is 1, but the data is 16-bit, we expect that each transfer actually involves double the number of physical columns. reg_packing: The number of elements that fit in a single 32-bit register. rrr rrN) rr'r(rrr-rrr rrr) base_addrcols atom_shape tmem_packing reg_packingr8 atom_rows atom_cols total_numregs_per_instr num_row_stepsmax_num lane_stepaddr_row num_processed instr_num remaining num_slice addr_row_cols r<_transfer_32xcolsr(s& ##B'##)Y Q  i)a=:a=0U__{5RS. i1   /- )&&(1, -' ~- .'' &izz)UWWi).C-JC%PQHMI!M1 19Q 6  !a   ! }y'@AiZZ EGGMI5EsKl )Y 99y m"M1 19Q 6 I %% % &sC"E%AE<EEc>tjjd}|jdk(r|jddk(sJ|jddz}dt j |jdjz}|dk(rd}tjd|jddft }tj|d}tj|d} tj|D]A\} } tj | ||g| d<tj | | |g| d<C|j#dd|jddj%dd}|jd dd k(sJ|dk(sJd } nX|dk(rHd }|j#dd|jdj%dd}d|cxkrdksJJ|dk(} n t'|t)||d||} | D]'\}}}}|||fj}t+||||| )y)Nrrrrr rrbrJrr)rrFrIrr)rr'r(ndimrrrflatrremptyobjectr r. ndenumerater extractelementrswapaxesr)r(rd)r vector_regsrr8rr store_shaper_c0c1rvregrcitr'r$r!r& regs_slices r<rr%s1 ##B'#   Q ;#4#4Q#71#<< <   1  !$enn[%5%5a%8%=%=>>+AK 88Q ))!,a0 ?D Q B Q B^^K06 T++D"5d9S9!9o++D"5d9S9!9o6 <<1k//2A 6 ? ?1 ED ::bc?f $$ $ 1   FaK   q![%6%6q%9 : C CAq ID ! !! !! ! Q F k **D'<M"79J3lIy)i*+00J k9j&IJr>ctjjd}|jdk(sJtj|j dj }|j\}tj|j}d|z}d|f}|dk(r|dk(rdgt|dzz} tj|d} tj|d} t|D]@\} } tj | | | d| z<tj | | | d| zdz<Bn#|Dcgc]}tj"||} }|dk(sJd}n|dk(rm|dk(sJ|Dcgc]}tj"||} }|dk(rd|cxkrdksJJ|dk(}n|dk(r|dk7rt%d|d||d|zk(sJd}nX||k7rt%d |d |d |tj|dk(sJ|Dcgc]}tj"||} }d}t| |z}t'|||||}|D]"\}}}}|dk(sJ| |}t)|d |||$ycc}wcc}wcc}w) Nrr rrFrrz for element type Only z packing supported for bitwidth , but got TMEM packing of rH)rr'r(r,rr-rrrrrlrr r. enumerater r1rr)r(rd)rr3rr8vec_tyrk elt_bitwidthrstore_atom_shaper_r5r6rr7rZrcrr8r'r$r!r&r9s r<rrJs ##B'#   Q   ==))!,11 2&LL/= 3 34,l"++&AVs;'!+ ,d >>#q !b >>#q !b -:)#t++D"5QW  //b9QWq[:.9 9emmAs# 9d 9 1   Fa A  +6 7aEMM!S ! 7D 7r , #! ## ## #q f  |r1!#L>1CL> R  R</ // /f{"  +>|nM&&2^ 5  >>& !R '' '+6 7aEMM!S ! 7D 7 F T[ $D*:L+V"79G3lIy) >>iJ h :vFG9: 8$ 8s.I&'I+I0ctjjd}tjj d|}dt j |z}|dk(r d}|dk(sJd}n(|dk(rd}d|cxkrdksJJ|dk(}n t|tjd|d zft } t||d ||} tj|d } tj|d} | D]\} }}}t| |||}t|dz|dzdz}| ||f}|j d|fk(sJ|j |f|dk(r|Dcgc]}t#j$||}}tj&|t j)|ddj+d d}t#j,|}|j g|j dk(sJtj.|j D]E}t#j0||g|d | }t#j0||g|d| }|||<GY|dk(sJ|Dcgc]}t#j$||}}tj&|t j)|dj+d d}||d <| Scc}wcc}w)Nrrr rJFrrIrrr*r+r.)rr'r(rr%rrr)rndarrayr/r(r r.r`rrr rasarrayrr2 mlir_undefr insertelement)rrrrr8r>r load_shaper\r3r8r5r6r'r$r!r&r_ row_slicevector_regs_updaterZundefr high_undefr7s r<rrs ##B'# ==  T5 )&ennU+++AJ 1   DaJ ! !! !! ! 1 D k ** Atqy>8+D'<M" ~~c1" ~~c1"79%3lIy) lJ 4 @Di!mi!mq%89I$Y %9:  # #9~ 5\8J8P8PR[7\\ 5a.2 3dll5!$ 3d 3ZZF + 3 3Iq! D M MaQR Sdoof%e ZZ9/559q9 99 9.445'#''tIcI1IC !!*d9S9!9orB"&3' A  /3 4!dll61% 4d 4ZZF + 3 3Iq A J J1a Pd $1%4 '45s J6J;cXtjjd}tjj |f|}dt j |z}||zdk(sJd}d|f} |dk(rd|cxkrdksJJ|dk(} n||k7rtd|d|d|d } t||| ||} tj|d} tj|d} dg||zz}| D]z\}}}}|dk(sJ|t|||| }|dk(r-|dk(r(|Dcgc]}tj||c}||<T|Dcgc]}t j||c}||<||dk(r|dk(rtj|dzft }tj"|}t%|j&D]D}tj(||d|z| }tj(||d|zdz| }|||<F|S||k(sJtj*|t }|Scc}wcc}w) NrrrHrr r;z supported for element type r<Fr*)rr'r(rr%rrr)r(r r.r`r rrrBr/rDrsizerErC)rrrrrkr8r>rrFload_atom_shaper\r8r5r6r_r'r$r!r& instr_regsrZr3rIrrJr7s r<rrsa ##B'# ==  m-u 5&ennU+++  $ )) )*%/A ! !! !! ! 1 D{"  +:5'B*^ -  DD/<U" ~~c1" ~~c1" 4;& '$79G3lIy) >$9$>L*iFJaMQ.9CDAeQ/Dd9o;EFaq&1Fd9o GA-1,**dai\8K OOF #E[%%&%%eT!c']B?j    DS1,=r Bdk#  K '' '**T0K EFs H"=H'ctjtjjt j yrf)r tcgen05_waitTcgen05WaitKindSTORErwarpgroup_barrierrr>r< commit_tmemrTs(D((../r>ctjtjjt j yrf)r rPrQLOADrrSrr>r<wait_load_tmemrWs(D((--.r>c tjjd}tj|j}|j x}|j k7rtd|d|j |tjjtjjhvrtd|d|jdtzr"tdtd|jd|jd d zrtd |jd |jtk7rtd |jd t!|j}|jdtz|jd d zddf}||k7rtd|d|d|j|j#\}} |t%j&|k7r td|dd\} } | dzs| dzr td| d z} | d z} t%j(|d}|ddkDr tdt+j,|ddD] \}}t%j.||| z|| zzg|}t1j2|j4t1j6|d |dz|zd |zz}t9j:|ddd}t=j>t<j@jBtE||t<jFjH|rt<jJjLnt<jJjN y)aAsynchronously copies the scale data from SMEM to TMEM. The result of the copy can be awaited by calling ``commit_arrive`` and waiting on the chosen ``Barrier``. However, if TMEM reference is to be consumed by a MMA issued in the same thread, no additional synchronization is needed. At the moment the function requires ``smem_ref`` to be contiguous and have a shape of ``(MN // 128, K // 128, 32, 16)`` for 8-bit scales (here MN stands for the size of the non-contracting dimension which is M or N), matching the scale layout for .scale_vec::1X. See https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-mma-scale-factor-a-layout-1x for more details. Note that we always put the non-contracting dimension first. If you have a (MN, K // 32) array of scales in JAX (where MN and K are divisible by 128), you can prepare it for use in the kernel this way:: scales.reshape(mn // 128, 4, 32, k // 4, 4) .transpose(0, 3, 2, 1, 4) .reshape(mn // 128, k // 4, 32, 16) The TMEM ref is expected to have the logical shape of the scales ``(MN, K // 32)``, and the layout created by ``scales_layout()``. rIncompatible dtypes: SMEM has , TMEM has Unsupported dtype: z+, only f8e8m0fnu and f8e4m3fn are supportedrz'TMEM reference must have a multiple of rr rz:TMEM reference must have a multiple of 4 columns, but got rrr SMEM has , but expected  for TMEM ref shape 9Only copies from contiguous SMEM references are supportedNrrz,Scale tile strides must be a multiple of 128rzOnly M/N up to 256 supportedrp) multicastr)(rr'r(rrrlrr-rCr%r+r)rrurrrget_strides_and_offsetrget_contiguous_stridesr7rr getelementptrr rrr.rencode_descriptorr tcgen05_cpTcgen05CpShape SHAPE_32x128br@Tcgen05CpMulticastWARPX4r%r&r8)smem_reftmem_refr_r8smem_tyr smem_shapeexpected_smem_shaperrXmn_tile_stride k_tile_stridemn_tile_stride_i32k_tile_stride_i32 smem_base_ptrmn_tilek_tileload_ptr store_addrr9s r<async_copy_scales_smem_to_tmemrys@0 ##B'# MM(-- ('###e6 5eWKGWX YY 2''++-r/B/B/F/F/HII  3E7:ef gg ^^A" >ykYaYgYghiYjXkl mm ^^A QRZR`R`abRcQde ff __ ' |HOO#44EF GGW]]#*!*i79Ja9OQSUWX&&  J</B.CDnn% ' --/*'1 ,,Z88 P QQ")"1+.-c]S0 C DD%*#q(""8Q/-]Q < ==JrN3ogv"" % %1B(B BC H  sA 1 -6WDEJ  & &xFD IDOO ))*% ))00)3d%%9J9J9P9P r>c *tjjd}tjjd}tj|j}|j x}|j k7rtd|d|j |tjjdk7rtd|d|jdd zrtd |jd|jd d zrtd |jd |jtk7rtd|jdt|j}|jdd z|jd d zd d f}||k7rtd|d|d|j|j\} } | tj|k7r td|dd k7r td| d } | dzr td| dz} tj |d} t#|d D]}tj$| || zg|}t'j(|j*t'j,|d|z}t/j0|dd d}t3|}t5j6t4j8j:|||rt4j<j>nt4j<j@y)NrrrYrZrr[z, only i2 supportedrrz9TMEM reference must have a multiple of 128 rows, but got r rbz:TMEM reference must have a multiple of 64 colums, but got rrr\r]r^r_zOnly M=128 supportedrz&K tile stride must be a multiple of 16rrr`r3)!rr'r(rrrlrr-r)rrrrrbrrcr7rrdr rrr.rrer@r rfrgSHAPE_128x128br%r&r8)rkrlr_rr8rmrrnrorrXrqk_tile_byte_stridertrvrw store_ptrr9r s r<'async_copy_sparse_metadata_smem_to_tmemr~6s ~~""1%" ##B'# MM(-- ('###e6 5eWKGWX YY bnn))!,,  3E7:MN OO ^^A PQYQ_Q_`aQbPcd ee ^^A QRZR`R`abRcQde ff __*,, |HOO#44EF GGW]]#*!*c18>>!3D3JCQST&&  J</B.CDnn% ' --/*'1 ,,Z88 P QQq 4 55!*-R = >>$)""8Q/-)!,- f""!334bH 8++U^^CV-LMI  & &xFD ID I &COO **C)3d%%9J9J9P9P r>)FFN) r/rr0rr3rr4rr5z int | Nonerr)rDzCallable[[ir.Type], int]r/rr0rrErrFrrGrrHrr3rr4rrIrrr)rr)rr`rzir.Value | TMEMRefrrrYrrZrr[TMEMRef | Noner\rr]rr^zir.Value | boolr_rrr)&rrrztuple[ir.Value, int] | ir.Valuertuple[ir.Value, int]rwrrxrryz.tuple[tuple[int, ...], tuple[int, ...]] | Nonerzz'tuple[tuple[int, ...], tuple[int, ...]]r{ir.Value | Noner|rr}rr/rr0rrvrrlrr~zir.Type | Nonerurr^rr_rrr)FN)r'zutils.BarrierRef | ir.Valuer_rr(zLaunchContext | Nonerr)r.rr/rrr)FT) r:rr.rr_rr/rrr)r:rrr) r:rr.rr_rr/rrr)r_rrr)rztuple[int, str])r\r)rcrrr)rrgrrrdrrr)rrr_rrdrrrg)r )rdrrrg)rrrdrrrg)rrrrr) rrrrrrrrrrrz*Iterator[tuple[ir.Value, int, int, slice]]r)rz np.ndarray)F)rkrrlr`r_rrr)K __future__r dataclassesrrr5typingrrrr jaxlib.mlirrjaxlib.mlir.dialectsr r r r numpyrrrrqrrlaunch_contextrru TMEM_MAX_COLSTCGEN05_SMEM_DESCRIPTOR_BITTCGEN05_LAYOUTrTCGEN05_TRANSPOSED_LAYOUTTRANSPOSED_LAYOUTTCGEN05_ROW_LAYOUT ROW_LAYOUTTCGEN05_COL_LAYOUT COL_LAYOUTrr=rKrTrXrrr+r0r<r@rCrFrUr`rd dataclassrrgrrrrrrrrr`r(rrrrrTrWryr~rr>r<rsr # 00&%'%$)   %  00  " "  " " **$(;? ;? ;?  ;?  ;?";?;?|(?0(? (? (?  (?  (?  (?(?(?(?(?(?VNN""(,"&CCCC  C  CCC&C CC CL e' e'3e' !e' e'  e' @ e'9e'"e'"e'#e' e' e' e'e''e' !e'"#e'$%e'& 'e'T $! (!! ! !6,P$* 4. J d#00$0f B B&) B47 B B* $$  "d#]H]H$]H@+&+& +& +& +&  +& 0 +&\"JJ2Gj.b)X  ?DMM")M7;M Mb?D,,"),7;, ,r>