ukit dZddlmZddlZddlmZddlZddlZddlm Z m Z m Z ddl m Zddl mZdd l mZdd l mZdd l mZe Zej,d GddZGddej0Zej,d GddeZej,d GddeZej,d GddeZej,d GddZej,d GddZej,d GddZej,d GddZ eezezezeze zZ! d;d Z" dd#Z%ej,d Gd$d%Z&e'ejPejRfejRejPfejTejVfejVejTfejXejPfejZejXfejZejPfgZ.ej,d Gd&d'Z/ej,d Gd(d)Z0ej,d Gd*d+Z1ej,d Gd,d-Z2e&e/ze1ze0ze2zZ3 d?d.Z4ej,Gd/d0Z5e Gd1d2Z6 d@d3Z7dAd4Z8 dBd5Z9 dCd6Z: dDd7Z;dEd8Z< dFd9Z= dBd:Z>y)Gz1Defines expressions and constraints over layouts.) annotationsN)Sequence)Any assert_neverfinal)fragmented_array)launch_context)layouts)inference_utils)tcgen05T)frozenc"eZdZUdZded<dZy)VariablezMA variable is an abstract identifier. `key` is supposed to be hashable. VariableKeykeyc"d|jdS)NzV()rselfs b/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/jax/experimental/mosaic/gpu/constraints.py__str__zVariable.__str__/s z N__name__ __module__ __qualname____doc____annotations__rrrrr'srrceZdZdZy)ConstantzA constant is a known layout.N)rrrrr!rrr#r#3s%rr#c"eZdZUdZded<dZy)RegisterLayoutzWraps a known register layout.fa.FragmentedLayoutvaluec"d|jdSNzC(rr'rs rrzRegisterLayout.__str__=  |1 rNrr!rrr%r%7s& rr%c"eZdZUdZded<dZy) TMEMLayoutzWraps a known TMEM layout.tcgen05.TMEMLayoutr'c"d|jdSr)r*rs rrzTMEMLayout.__str__Gr+rNrr!rrr-r-As" rr-c"eZdZUdZded<dZy) SMEMTilingzWraps a known SMEM Tile Transform. If an SMEM reference may, in principle, have transforms but should not be tiled, then `value` is `None`. lc.TileTransform | Noner'c"d|jdSr)r*rs rrzSMEMTiling.__str__Ur+rNrr!rrr1r1Ks ! rr1c(eZdZUded<ded<dZy)Reduce Expression expressiontuple[int, ...]axesc<d|jd|jdS)NzReduce([z], r)r9r7rs rrzReduce.__str__^s dii[DOO#4A 66rNrrrr rr!rrr5r5Ys7rr5c,eZdZUded<ded<ded<y)BroadcastInDimr6r7r8r9shapeNrrrr r!rrr=r=bs rr=c,eZdZUded<ded<ded<y)Reshaper6r7r8 source_shape target_shapeNr?r!rrrArAisrrAceZdZUded<dZy) Transposer6r7c"d|jdS)NzT(rr7rs rrzTranspose.__str__ts   ""rNr;r!rrrErEps #rrEcdfd }tj|}|xtdx \tStdxy\}|tj dx>\}||s tStt j j S t|jj S t|jj S)Ncntj|dD]\}}j||k7syy)NTstrictF)zipr9r>)r>axiss broadcasts r_check_shape_broadcastz;reduce_broadcast_expression.._check_shape_broadcasts<y~~uT:a  ! # rr!r*r>)r7r9r>)r>r8returnbool) reduce_expressionr7 Unsatisfiabler%faWGSplatFragLayoutr>r=r9)rO assignmentsrP reduced_exprlayoutr>s` rreduce_broadcast_expressionr[s #9#7#7E,  _   %   ! ! .'. ? " 4 49?? KL L/%>>OO   & ! ioorct|j|}|xtdx \tStdx\}|xtj dxd\}t j|t j|jk(sJtt j |jSxtjdxf\}}t j|t j|jk(sJtt j|j|Sxtjdx\}|j}t|jt|krtj||St||j t| ddD]&\}} | |zdk7stj||cSt|d z } |j | d} |j| d} |d} | | k7s|jt| | zdk7rtj||St|Sy tj||S) Nr!r*rQ)r>vec_sizerGTrJrr)rTr7rUr%rVrWmathprodrCWGStridedFragLayout TiledLayoutbase_tile_shapelen dataclassesreplacerLrB)reshaperXrYrZr>r] tiled_layout tile_shapetsrNnum_minor_tiled_dimssource_minor_tiled_dimstarget_minor_tiled_dimsmajor_tiled_dims rreduce_reshape_expressionrnsE#7#5#5{C,  _   %  .R ! ! .5!TYYw/C/C%DD DD"")=)=>  / DR # # C C5!TYYw/C/C%DD DD$$,,x  D.R^^  -#33* %% &Z 8&&w<H H :w';';S_ra layouts_libsplat_is_compatible_with_tiled_SUPPORTED_TILED_RELAYOUTS)rrr source_layout target_layoutsplatstrideds rrzRelayout.holds2s [[F [[F  fn -Z6#)<<=M  & O O *2   ! *,OB,B,B,D O{{gmm++ P 3 32   !2>>#3 399 =   4 . -2>> R^^-}-1KKK . rc<d|jd|jdS)Nz Relayout( ⟶ rrrrs rrzRelayout.__str__Ps t{{m6$++a 88rNrrr!rrrr"s   <9rrcveZdZUdZded<ded<ded< d dZ ddZ dd Zdd Zd Z y )IsTransferablezYStates that `source` layout must be transferable across memory spaces to `target` layout.r6rrr8r>cdt|jdk(sJ|jd}tj|}|tj f|tj ftj||tjftj||tj|fgS)z:Returns the list of supported TMEM <-> Register transfers.rpr) rcr>r tmem_default_layoutrVTCGEN05_LAYOUTTMEM_NATIVE_LAYOUTtmem_half_lane_layout WGMMA_LAYOUTtmem_m64_collective_layoutfa_m64_collective_layout)rpackingcolumnsrs rsupported_tmem_transfersz'IsTransferable.supported_tmem_transfers]s tzz?a  jjmG!55g> b//0 b334  & &w 8"//J  . .w @  , ,W 5  rcD|j}||f|j|vSN) vector_lengthr)r tmem_layout reg_layoutrs r_is_valid_tmem_transferz&IsTransferable._is_valid_tmem_transferns*''G  $(E(Eg(N NNrcptj|r|duxrt|jdk(S|duS)Nrp)r is_mma_layoutrcrr)r smem_layoutrs r_is_valid_smem_transferz&IsTransferable._is_valid_smem_transferts=$$Z0  $ E[-?-?)@A)EE $ rc|j|jk7sJd|j|jfx r4dk(r/\tdx#\ctdx\}}|j ||Sx r4dk(r/\tdx#\ctdx\}}|j ||Sx r4dk(r/\t dx#\ctdx\}}|j ||Sx r4dk(r/\tdx#\ct dx\}}|j ||S rmdk(rh\tdx\\tdxR\t|jj}t|jj}td|d| y)zbReturns whether the constraint holds. Returns `None` if the constraint can't be checked. zJIsTransferable constraints within the same memory space are not supported.rpr*Nr!zUnsupported transfer: z -> ) rrr-r%rr1rr#typerrs)rsrcdst source_type target_types rrzIsTransferable.holdssL ;;$++ %  % ++t{{ " ; ;: C."; ;++C55 < ; ;> $j&; ;++C55 < ; ;: C."; ;++C55 < ; ;> $j&; ;++C55 < ! !8:xz4;;'00 4;;'00 !$[Mk] C   " rc<d|jd|jdS)NzIsTransferable(rrrrs rrzIsTransferable.__str__s T[[M }A >>rN)rintrRz4list[tuple[tcgen05.TMEMLayout, fa.FragmentedLayout]])rr.rr&rRrS)rr2rr&rRrSr) rrrrr rrrrrr!rrrrTsxa   ;"O+O9LO O * &  :?rrc4eZdZUdZded<ded<d dZdZy) NotOfTypez0States that `expr` is not an instance of `type`.r6r}ztype[fa.FragmentedLayout]rct|jtsyt|jtsyt|jj|j  S)zjWhether the distinctiveness constraint holds. Returns `None` if the constraint can't be checked. NT)r{r}r#r%r'rrs rrzNotOfType.holdssB dii *  dii 0 $))//4995 55rcNd|jd|jjS)Nztype(u) ≠ )r}rrrs rrzNotOfType.__str__s$ 499+VDII$6$6#7 88rNrrr!rrrrs8!! 69rrc4eZdZUdZded<ded<d dZdZy) DividesaStates that the `expr` tiling is a divisor of `tiling_multiple`. That is to say that, for each tiled dimension in `expr`, the dimension must divide its corresponding dimension in `tiling_multiple` starting from the tail. If `tiling_multiple` contains more dimensions than `expr`, then the extra dimensions in `tiling_multiple` are ignored for the purposes of the check. `expr` is not allowed to contain more dimensions than `tiling_multiple`, and this constraint therefore also constrains the rank of `expr`. r6r}r8tiling_multiplec|jxtdx\yxtdx\tjdx\}|}nPxtdx(\xt j dx\}|j}ntdx\}|j}n yt|t|jkDrytt|t|jD] \}}||zs yy)Nr*T)rrr!F) r}r1rtrur%rVrarbr-rcrrLreversed)rtrrrZr'sizemultiples rrz Divides.holdss )) !: !D ! " 8: 7B,,6 7 8 ;> ; : 0 ;'' <  &&   6{S--..hv.9M9M0NOh D rc:|jd|jdS)Nz % z == 0)rr}rs rrzDivides.__str__s "" #3tyyk 77rNrrr!rrrrs  ""08rrcr|xtdx]\}}t||}t|tr tSt||}t|tr tSt||SxtdxS\}}t||}t||} t|tst| tr tSt || Sxt dx7\} } t| |} t| tr tSt | | Sxt dxU\}}} t||}t||} t|tst| tr tSt || | Sxtdx7\} }t| |} t| tr tSt| |S}t|y)zReduces a constraint.rNrr}rrrr>r}r) rrTr{rUrrrrr) constraintrXrrlhs_redrhs_redrr source_red target_redr}rexpr_redr>rnevers rreduce_constraintrs  ! ! !!#{3g G] +!#{3g G] + GW %% " 0 / /$V[9j$V[9j J .* m3 j* -- 0 ) ( ("45h Hm , x && ) C B B$V[9j$V[9j J .*Z2W J E :: C = < <"45h Hm , X // = 5rceZdZUdZej eZded<ej e Z ded<d dZ d dZ d Z y ) ConstraintSystema.A constraint system contains a set of constraints and assignments. Assignments assign constant values to variables in the system (bound variables). Constraints describe relationships between variables that must be upheld, and can be used to determine assignments for unknown (free) variables. )default_factorydict[Variable, Constant]rXSequence[Constraint] constraintsc t g d fd jD]}|xtdx\}}||$xtdx\}}||Extdx \}|]xt dx\}}||xt dx \}|}t| S)z1Returns the list of free variables in the system.c|xtdx:\|vr2|jvr#j|j|yyyxtdx\yxt dx \}|yxt dx \}|yxtdx \}|ytdx \}|y t|y)Nr!rG) rrXaddappendr#r5r=rArEr)r}eextract_variablesfree_variablesseen_variablesrs rrz4ConstraintSystem.unknowns..extract_variables's  XZ  'D8H8H,H   t $  ! !$ '-I 'XZ  !V ! A " )^ ) A * "W " A #  $ A %  t rrr)r}r)r}r6rRNone)setrrrrrrr) rrrrrrr}rrrrs ` @@@runknownszConstraintSystem.unknowns#s$'EN%'N$&&  %V % % C C &4X 3 3 F # F #4"Y ! D !" B^ B B F # F #C W  D !  u !" rc|jjD]2\}}||jvs||j|k7s(tcSt|j|jzg|j|jS)NrXr)rXitemsrUrr)rothervariable assignments r__and__zConstraintSystem.__and__Ls!% 0 0 6 6 8* U&& &:9J9J89T+T $$u'8'88;d&&;):):; rcd}|dz }|jjD]\}}|d|d|dz }|dz }|jD] }|d|dz } |S)NzConstraintSystem z assignments: z u ⟵  z constraints: )rXrr)rrrconstantrs rrzConstraintSystem.__str__WsA A $ 0 0 6 6 80 HT*U8*B //a0 A&&! T*R  a! HrN)rRzlist[Variable])rrrR ConstraintSystem | Unsatisfiable)rrrrrdfielddictrXr listrrrrr!rrrrs`+<+*;*;++''8k&7&7&M+#M'R # '  rrceZdZddZy)rUc|Srr!)rrs rrzUnsatisfiable.__and__es KrN)rrrRrU)rrrrr!rrrUrUbsrrUct}|D]U}|tdxH\xtdx:\ctjk(r'}t |tsJ|j |SW|S)z)Returns a all vars distinct from a splat.rr!)rrrrVrWr{r)rvarsrvars rnon_splat_variablesris`$j   G+(*+"2F2F G#x(((   H +rcvt|}|sydd}|D]"}|tdx\}}||vs||sy $y)a Returns whether the constraints imply a non-splat to splat relayout. Such relayouts are impossible and this helps shortcut the search. If this function returns False, this doesn't necessarily mean that there are no non-splat to splat relayouts, just that this is not known yet. Fcnt|txr$t|jtjSr)r{r%r'rVrW)rs ris_constant_splatz>_has_relayout_of_non_splat_to_splat..is_constant_splats, a ( Z %%.rrT)rRrS)rr)r non_splatrrrrs r#_has_relayout_of_non_splat_to_splatrvsX"+.)     j   1 Y #4V#< 2    rcjt|j}g}t|dkD}|r~d}|jD]j}|tdx\\}}t |t s"||vs'||vs,d}|j ||jt|tji l|r~|t|zS)aAdds transitive NotOfType constraints for all non-splat variables. Given `n` variables `l0`, ... `l{n-1}`, and a set of relayouts `{ Relayout(l{i}, l{i+1}) : 0 <= i < n }`, if we also know that `l{0}` is not splat, then we can automatically deduce that none of `l0`, ..., `l{n-1}` are splat either. This helps us quickly conclude that a system is unsatisfiable in cases where a non-splat variable is transitively relaid out into a splat layout. rFrTr) rrrcrr{rrrrrVrWr)constraint_systemrnew_constraintsnew_non_splat_foundrrrs rsaturate_distinct_from_splatrs""3"?"?@)&(/I*'33    3*I% )"&  MM& !  " "9VR5I5I#J K4    -/J JJrcN i d fd d fd }t}|jD]|}|tdxo\xtdxa\cxtdxT\}}t |tsJt |tsJ|j ||j ||||z~i}t |tD]+} |}|j|gj|-i} |jD]"} | D]}| D cgc] } | |k7s |  c} | |<$| Scc} w)zComputes all transitively equal variables in a constraint system. The output dictionary maps each variable that appears in constraints in the constraint system to all the variables it is transitively equal to. cL|vr||<||k7r||<|Srr!)vfindparents rrz-compute_transitively_equal_vars..finds9fQi ayA~vay/fQi !9rc<|}|}||k7r||<yyrr!)v1v2root1root2rrs runionz.compute_transitively_equal_vars..unions* HE HE ~fUmrrr!r)rrrRr)rrrr) rrrrr{rsortedstr setdefaultrvalues)systemr all_varsrrr componentsrroot equal_varscomponent_varsrrrs @@rcompute_transitively_equal_varsrs1&(&  E(&&j   ?'hj'->XZ ?#x(((#x((( S S c3 @02* ( $.a 7D$#**1-.02*"))+Hn H*8GEQJuGjmHH Hs  D"D"cLt|}g}|jD]g}|j||tdxH\}}t |t s3|j |gD]}|jt||f it|}tj||S)zESaturates Divides constraints between all transitively equal vars. rr) rrrrr{rrzmerge_divides_constraintsrdre)rrrrr}r equal_vars r+saturate_divides_constraints_for_equal_varsrs /v6*&(/&& j:&   > dH %%>>$3 Hi  " "79o#F G H ?   .o>/   V AArc$g}i}|D]}|tdx\xtdx\}}t|tsJ|j|x}|||<Ht t |t |}g}|dkDrCt || d|| ddD]*\} } |jtj| | ,t|||< |j||jD] \} }|jt| |"|S)z.Merges Divides constraints that can be merged.rNr!rTrJ) rrr{rzminrcrLrr^gcdtupler) rresultvar_to_tiling_multiplesrrrprevious_tiling_multiplemin_lennew_tiling_multiplexyr}s rrrs1&>@"j   I I!X&&&(?(C(CA(F F $ O'6 !! $  c/*C0H,IJ  Q;/7()46NPWxy6Ybfg 7da  & &txx1~ 6 7%*+>%?" J  j!%"& 7<<>2dO MM'$012 -rc  |j g}d}d fd }|jD]}t| xtdx\tcSxtdx8\xt dx*\cxt dx\}}|||s tcSd}ixtdx8\xt dx*\cxt dx\}}|||s tcSd}}t|tsJ|jx|j||||k7z}xdur tcSdusd}t|}|t|t|k7z}|}t|r tS|rt |jz|Sy) a6Performs one reduction step over each constraint in a constraint system. Returns: - Unsatisfiable(): if the constraint system is unsatisfiable. - A new constraint system if any constraint was reduced. - None: if the constraint system is not known unsatisfiable, but hasn't been reduced. Fc*|vr ||k7ry||<y)NFTr!)rcstrXs r try_assignz'_reduce_system_once..try_assign's' kk#.#5 K rr!NrTr)rrr'r#rRrS)rXrrrUrrr#r{ Constraintrrrrcrr) rrchangedr(rrr'new_constraintrrXs @r_reduce_system_oncer,s"--+"$+ ' &11j J 4 =?  ?6 ?'hj'->XZ ?#s#  @ @6 ?'hj'->XZ ?#s#  @ .*555""$   ~ . ~3 3G ? "G+..k:/ S !S%5 55'+)5 ? "3"?"??  rc t|x |Sxtdx \tSxxtdx\}|}n}t|I)zReduces a constraint system until it can no longer be reduced. Returns: - Unsatisfiable(): if the constraint system is unsatisfiable. - The maximally reduced constraint system otherwise. r!)r,rUrr)r new_systemrs rr|r|Us[ / 0    =?  +   +& , U r)rOr=rXrrRExpression | Unsatisfiable)rfrArXrrRr/)rvrErXrrRr/)r}r6rXrrRr/)rr)rXrrRzConstraint | Unsatisfiable)rrrRz set[Variable])rrrRrS)rrrRr)rrrRzdict[Variable, list[Variable]])rrrRr)rrrRzlist[Constraint])rrrRz'ConstraintSystem | Unsatisfiable | None)?r __future__rabccollections.abcrrdr^typingrrrr rVr rtr rr r r dataclassrABCr#r%r-r1r5r=rArEr6r[rnrxrTr frozensetrWGMMA_TRANSPOSED_LAYOUTrTCGEN05_TRANSPOSED_LAYOUTWGMMA_LAYOUT_UPCAST_2XWGMMA_LAYOUT_UPCAST_4Xrrrrrr)rrrUrrrrrrr,r|r!rrr<s8 # $ ++$"$ d#$&sww&d#X$d#$d#  $ d#77$7d#$ d#  $ d###$#       ,D>2C 2C#;2C2Cj00'?00( #;6d# / /$ /'__b0011445!!2#4#450 9 9:0 ( d#.9.9$.9bd#L?L?$L?^d#99$9*d#*8*8$*8Zh  *^ ;g E '')A''T I I I X %  6 K' K% KF0 0#0fB BB*8:':,:z'%r