uki~4ddlZddlmZddlmZddlmZddlmZddl Z ddl m Z Gdd Z d ejd ejd ejd ejfdZd ejd ejd ejd ejfdZy)N)fragmented_array)ir)llvm)vector)utilsceZdZdZej ej ddddZej ej dejdfddZ ej ej d dddZ y ) MMALayoutszzContainer for MMA layouts, providing a convenient way to create layouts for MMA operands based on warp configuration. )@r rr)i)) warp_dims lane_dims vector_dim)rr rr)r rrrrN) __name__ __module__ __qualname____doc__fa TiledLayoutTilinglhs ReplicatedrhsaccZ/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/experimental/mosaic/gpu/mma.pyr r s bii12  # bii'(q!#  # bii01  #r(r r&abreturnc : |jdk(sJ|jdk(sJjdk(sJ|j|jk(sJ|jtjj tj j fvsJjtj j k(sJtjtjrHt|jtjr$t|jtjsJd\}}}jjDcgc]\}tjjD]8}tj |gtj"j |g:^}}}tj$j'd} |jjD cgc]} t)j*| | } } |jjD cgc]} t)j*| | } } t-| dk(sJt-|dk(sJt-| dk(sJd |jd |jd } t/j0fd }||}||}||}||}| d |d|d|d|d }dj3dg|zddj3dg|zddj3dg|zddj3dg|z}g| | |}tj4j7ddj3fd|Dd}t9j:||||d}tt-|Dcgc]$}t9j<j||g&}}g}t9j>tj@j dj}tC|ddd|dddD]n\}}t9jD||t)jFd| }t9jD||t)jFd| }|jI|ptKjL|tNjQjj}tjR|jdScc}}wcc} wcc} wcc}w) z;Performs `acc + a @ b.T` using warp level MMA instructions.r rr)rrr)dynamic_positionstatic_position rrz&mma.sync.aligned.m16n8k16.row.col.f32..z.f32c ddjt|Dcgc]}dtc}zdzScc}w)N{,$})joinrangenext)n_counters r)z"_mma_single_tile..Ys; CHH58z#_mma_single_tile..ksFCNN 3Fs"z)>F)has_side_effectsrNrr)position)dtype) _registers_layout _is_signed)*shaperErF16TypegetBF16TypeF32Type isinstancelayoutr r! registersflattenr8 vector_lengthrextractDenseI64ArrayAttr IntegerType get_signlessrbitcastlen itertoolscountr7Typeparser inline_asm extractvalue mlir_undef VectorTypezip insertelementcappendnpasarrayobjectreshapeFragmentedArray) r&r*r+ num_acc_regs num_a_regs num_b_regsregposacc_regsi32rAa_regsb_regsinstr n_regs_str out_regs_str a_regs_str b_regs_str c_regs_strptx constraints in_operandsacc_struct_typeout_regs_structiout_regsvec_regs vec_undeffirstsecondvecr<s ` @r)_mma_single_tiler2s H   G   g    %% % "**..*BKKOO,=> >> > 2::>>+ ++ +R^^, QXXr~~ . QXXr~~ ./*1&, J&&(  szz//0   nn ..22C59( ##B'#+,;;+>+>+@ AaEMM!S ! A& A+,;;+>+>+@ AaEMM!S ! A& A V   X!   V   21<<.!,,t T% OO '*L),*%**%*,'* '<.:,b Bzl!L# 4&% &'q 3% " #$A 3% " #$A 3% $ %&(.&-6-H-+GGMMsxxFXFFGrJ/OO  /S]#  !=((oobmm//cnnEF)8CaC=(14a4.9meV   Y3 HC   S&5771c? CC OOCZZ / 7 7 8K8K L(  3::$ u B AHsA!T8T3T5)Tc |j\}}|j\}}|j\}}||k7rtd|d|||k7rtd|d|||k7rtd|d|tjj } tj j } |j |j k7r%td|j d|j |j | | fvr td|j tjj k7r tdtj|jk7r tdtj|jk7r td tj|jk7r td t|jtj sJt|jtj sJt|jtj sJ|jj"\} } |jj"\} }|jj"\}}| |k(sJ|| k(sJ|| k(sJ|| z|| z|| z}}}||k7rtd|d|||k7rtd|d|||k7rtd|d|| d k(r | d k(r| d k(sJd| d| d| d|j%}d}t'|D]^}t'|D]N}t'|D]>}||| }||| }||| }t)|||f|||f|||f|||f<@P`|S)a`Computes `acc + a @ b.T` using synchronouse MMA instructions. All operands must have `TiledLayout`s. The layouts must be generated by the `MMALayouts` class, which ensures that the tiles are mapped to the warps correctly. Args: acc: A `FragmentedArray` with a `TiledLayout` generated from `MMALayouts.acc`. a: A `FragmentedArray` with a `TiledLayout` generated from `MMALayouts.lhs`. b: A `FragmentedArray` with a `TiledLayout` generated from `MMALayouts.rhs`. Returns: A new `FragmentedArray` with the result of the computation with the same type as `acc`. z M mismatch: z != z N mismatch: z K mismatch: zDtype mismatch: z-Only bf16 and f16 supported for the operands.zOnly f32 accumulator supported.z$Expected MMALayouts.lhs layout for Az$Expected MMALayouts.rhs layout for Bz&Expected MMALayouts.acc layout for accr rr z Tile shape r?z not supported.c,t||z|dz|zS)Nr)slice)idxlengths r)r=zmma..s%f sQw&.@Ar()rN ValueErrorrrQrPrOrENotImplementedErrorrRr r#rTr%r&rSr r!base_tile_shapecopyr8r)r&r*r+mkr:k2m2n2bf16f16m_tilek_tilen_tilek_tile2m_tile2n_tile2 num_m_tiles num_n_tiles num_k_tilessk_idxm_idxn_idxmsnskss r)mmarsu. 77&1a GG'1b YY(2r"W |A3d2$/ 00"W |A3d2$/ 00"W |A3d2$/ 00  $ #\\Q\\! ' ~T!,,H II\\$$ M NN^^rzz~~'' ? @@^^qxx ; <<^^qxx ; <<^^szz! = >> AHHbnn -- - AHHbnn -- - CJJ // /88++.&&HH,,/&'ZZ//'7 7   F   F  *+v+qF{AKK{+"W |A3d2$/ 00"W |A3d2$/ 00"W |A3d2$/ 00 2&A+&B,F82fXRx? 6  #A![!Je{#J%J% uf  uf  uf &s2r6{Ab"fIqRyIBF JJJ *r()r^jax.experimental.mosaic.gpurr jaxlib.mlirrjaxlib.mlir.dialectsrrnumpyrjrr rnrrr'r(r)rs >%'2O  O " 2 2O797I7IOOfW  W  W  W  W r(