PiXddlmZddlmZmZddlmZmZmZddl m Z dZ GddZ ed Gd d e Z ed Gd d e Zed Gdde Zed Gdde Zed Gdde ZeddGdde ZGddZe dZed GddeZeddGddeZeddGddeZd#d Zd!Zd"S)$) dataclass)ListOptional)_unwrap_if_constexpr _unwrap_shapeconstexpr_type)constexpr_functioncB|jpdg|z}|jpdg|z}|jp(tt t |}t |d|t |d|t |d|dS)N ctas_per_cga cta_split_num cta_order)r r rlistreversedrangeobject __setattr__)layoutrankr r rs /var/www/development/aibuddy-work/election-extract/venv/lib/python3.11/site-packages/triton/experimental/gluon/language/_layouts.py_realize_cta_layoutrs&41#*L(6QC$JM ?D%++)>)>$?$?I v~|<<< v >>> v{I66666c(eZdZdZedZdS)DistributedLayoutz@ Base class for distributed memory layouts in Gluon IR. c t|SNrselfs rtypezDistributedLayout.typed###rN__name__ __module__ __qualname____doc__propertyr rrrr9$$X$$$rrT)frozenceZdZdZdZdS) AutoLayoutc*|Sr)get_auto_layoutrbuilders r_to_irzAutoLayout._to_irs&&(((rcdS)NALr(rs rmanglezAutoLayout.mangle strN)r#r$r%r1r4r(rrr,r,s2)))rr,ceZdZUdZeeed<eeed<eeed<eeed<dZeeeed<dZ eeeed<dZ eeeed <fd Z d Z d e fd ZdZxZS) BlockedLayouta` Represents a blocked layout, partitioning a tensor across threads, warps, and CTAs. Args: size_per_thread (List[int]): Number of elements per thread per dimension. threads_per_warp (List[int]): Number of threads per warp per dimension. warps_per_cta (List[int]): Number of warps per CTA per dimension. order (List[int]): The ordering of dimensions for partitioning. ctas_per_cga (Optional[List[int]]): CTAs per CGA grouping. cta_split_num (Optional[List[int]]): Split factors for CTAs. cta_order (Optional[List[int]]): Ordering for CTAs. size_per_threadthreads_per_warp warps_per_ctaorderNr r rc`tdt|jtdt|jtdt|jtdt|jtdt|jtdt|jtdt|j t|j}t||t|j|ksJt|j|ksJt|j|ksJt|j|ksJt|j|ksJt|j |ksJdS)Nr7r8r9r:r r r) superrrr7r8r9r:r r rlenrrr __class__s r __post_init__zBlockedLayout.__post_init__:s -/CDDX/Y/YZZZ .0DTEZ0[0[\\\ O-A$BT-U-UVVV G%9$*%E%EFFF N,@AR,S,STTT O-A$BT-U-UVVV K)=dn)M)MNNN4'((D$'''4())T11114%&&$....4:$&&&&4$%%----4%&&$....4>""d******rc ~||j|j|j|j|j|j|jSr)get_blocked_layoutr7r8r9r:r r rr/s rr1zBlockedLayout._to_irLsB))   !   J     N   rreturncd}||j}||j}||j}||j}||j}||j}||j}d|d|d|d|d|d|d|dS)NcZ|dSdtt|SN_joinmapstrxs r stringifyz'BlockedLayout.mangle..stringifyY&yr88CQKK(( (rB)r7r8r9r:r r r) rrOr7r8r9r:r r rs rr4zBlockedLayout.mangleWs ) ) ) $)D$899$9T%:;;! $"455  $*%% y!233 ! $"455 Idn-- {?{{%5{{ {{{{P\{{_l{{ox{{{{rc `tt|jt|jt|jt|j|jrt|jnd|jrt|jnd|jrt|jndfSr) hashtupler7r8r9r:r r rrs r__hash__zBlockedLayout.__hash__gs $& ' ' $' ( ( $$ % % $*  (,(9 CE$# $ $ $t)-); EE$$ % % %%)^ =E$. ! ! !    rr#r$r%r&rint__annotations__r rr rr@r1rLr4rU __classcell__r?s@rr6r6$s  #Y3i9 9(,L(49%,,,)-M8DI&---%)IxS ")))+++++$    |||||        rr6cPeZdZUdZeed<eed<fdZdZde fdZ dZ xZ S) SliceLayoutz Represents a layout corresponding to slicing a distributed tensor along one dimension. Args: dim (int): The dimension index to slice. parent (DistributedLayout): The parent layout before slicing. dimparentctdt|jtdt|jdS)Nr]r^)r<rrr]r^rr?s rr@zSliceLayout.__post_init__sS E#7#A#ABBB H&:4;&G&GHHHHHrch||j|j|Sr)get_slice_layoutr]r^r1r/s rr1zSliceLayout._to_irs2'' H K  w ' '   rrCcLd|jd|jdS)NSLrH)r]r^r4rs rr4zSliceLayout.mangles+7DH77t{11337777rc8t|j|jfSr)rSr]r^rs rrUzSliceLayout.__hash__sTXt{+,,,r r#r$r%r&rWrXrr@r1rLr4rUrYrZs@rr\r\ss HHH IIIII   88888-------rr\ceZdZUdZeeeed<eeeed<eeeed<eeeed<eeed<fdZdZd Z d Z xZ S) DistributedLinearLayouta Represents a linear distributed layout with explicit bases at register, lane, warp, and block levels. See: https://arxiv.org/abs/2505.23819 for reference. Args: reg_bases (List[List[int]]): Bases for register-level distribution. lane_bases (List[List[int]]): Bases for lane-level distribution. warp_bases (List[List[int]]): Bases for warp-level distribution. block_bases (List[List[int]]): Bases for block-level distribution. shape (List[int]): The tensor global shape. reg_bases lane_bases warp_bases block_basesshapec0tdt|jtdt|jtdt|jtdt|jtdt|jt|j}|jD]}t||ksJ|jD]}t||ksJ|jD]}t||ksJ|jD]}t||ksJdS)Nrirjrkrlrm) r<rrrirjrkrlrmr=)rrbasisr?s rr@z%DistributedLinearLayout.__post_init__s| Kt~)F)FGGG L-*H*HIII L-*H*HIII M=9I+J+JKKK G]4:%>%>???4:^ & &Eu::%%%%%_ & &Eu::%%%%%_ & &Eu::%%%%%% & &Eu::%%%%% & &rcf||j|j|j|j|jSr)get_distributed_linear_layoutrirjrkrlrmr/s rr1zDistributedLinearLayout._to_irs644T^T_VZVegkgw59ZAA Arc Xd|jd|jd|jd|jd|jd S)NDLLrH)rirjrkrlrmrs rr4zDistributedLinearLayout.mangles@lT^lldollll$JZll]a]gllllrc vtttt|jttt|jttt|jttt|jt|jfSr)rSrTrKrirjrkrlrmrs rrUz DistributedLinearLayout.__hash__s #eT^,, - - #eT_-- . . #eT_-- . . #eT-.. / / $*       r) r#r$r%r&rrWrXr@r1r4rUrYrZs@rrhrhs  DIT#YT#Yd3i   9&&&&&$AAAmmm       rrhcZeZdZUdZeed<eed<eed<fdZdZde fdZ d Z xZ S) DotOperandLayouta  Represents a layout for a dot operand. Args: operand_index (int): 0 for LHS and 1 for RHS of the dot operation. parent (DistributedLayout): The parent layout, representing the MMA. k_width (int): Number of elements per 32-bits. operand_indexr^k_widthc@tdt|jtdt|jtdt|jdS)Nrwr^rx)r<rrrwr^rxr`s rr@zDotOperandLayout.__post_init__sy O-A$BT-U-UVVV H&:4;&G&GHHH I';DL'I'IJJJJJrct||j|j||jSr)get_dot_operand_layoutrwr^r1rxr/s rr1zDotOperandLayout._to_irs2--d.@$+BTBTU\B]B]_c_klllrrCc\d|jd|jd|jdS)NDOrH)rwr^r4rxrs rr4zDotOperandLayout.mangles6PD&PP););)=)=PP PPPPrcDt|j|j|jfSr)rSrwr^rxrs rrUzDotOperandLayout.__hash__sT'dlCDDDrrfrZs@rrvrvs  LLLKKKKK mmmQQQQQEEEEEEErrv)r*eqceZdZUdZeeed<eeed<eeed<dZeeeed<dZ eeeed<dZ eeeed<fd Z d Z d e fd Zd ZxZS)NVMMADistributedLayouta Represents a layout for NVIDIA MMA (tensor core) operations. Args: version (List[int]): Version identifier for the MMA instruction. warps_per_cta (List[int]): Number of warps per CTA. instr_shape (List[int]): Instruction shape for MMA. ctas_per_cga (Optional[List[int]]): CTAs per CGA grouping. cta_split_num (Optional[List[int]]): Split factors for CTAs. cta_order (Optional[List[int]]): CTA ordering. versionr9 instr_shapeNr r rc\tdt|jtdt|jtdt|jtdt|jtdt|jtdt|jt|j}t||t|j|ksJt|j|ksJt|j|ksJdS)Nrr9rr r r) r<rrrr9rr r rr=rr>s rr@z$NVMMADistributedLayout.__post_init__s^ I';DL'I'IJJJ O-A$BT-U-UVVV M+?@P+Q+QRRR N,@AR,S,STTT O-A$BT-U-UVVV K)=dn)M)MNNN4%&&D$'''4$%%----4%&&$....4>""d******rcr||j|j|j|j|j|jSr)get_mma_layoutrr9r r rrr/s rr1zNVMMADistributedLayout._to_irs;%%dlD4FHY[_[m&*nd6FHH HrrCc hd|jd|jd|jd|jd|jd|jd S)NMMA_rH_MMA)rr9rr r rrs rr4zNVMMADistributedLayout.mangle sKdlKKT%7KK$:JKKTM^KKaeasKKvzwEKKK Krc :tt|jt|jt|j|jrt|jnd|jrt|jnd|jrt|jndfSr)rSrTrr9rr r rrs rrUzNVMMADistributedLayout.__hash__ sU4<((%0B*C*C4+,,$J[.eeD4E.F.F.Fae262DNU4-...$.2nFU4>***$HII IrrVrZs@rrrs   #Y9c(,L(49%,,,)-M8DI&---%)IxS "))) + + + + +HHHKKKKKIIIIIIIrrc(eZdZdZedZdS) SharedLayoutz; Base class for shared memory layouts in Gluon IR. c t|Srrrs rr zSharedLayout.typer!rNr"r(rrrrr)rrc|}|Wt|t|ksJtt|D]}||xx||zcc<|Sr)r=r)rmr shape_per_ctar]s r_get_shape_per_ctarsuM =!!SZZ////]++,, 5 5C #   -"4 4     rceZdZUdZeed<eed<eed<dZeed<dZeed<dZ e e eed <dZ e e eed <dZ e e eed <fd Zd Zee ddZdefdZdZxZS)NVMMASharedLayouta4 Represents a layout for shared memory suitable for NVIDIA MMA operations. Args: swizzle_byte_width (int): Width in bytes for swizzling. element_bitwidth (int): Bitwidth of element type. rank (int): Rank of the tensor. transposed (bool): Whether the layout is transposed. fp4_padded (bool): Whether FP4 padding is used. ctas_per_cga (Optional[List[int]]): CTAs per CGA grouping. cta_split_num (Optional[List[int]]): Split factors for CTAs. cta_order (Optional[List[int]]): CTA ordering. swizzle_byte_widthelement_bitwidthrF transposed fp4_paddedNr r rc>tdt|jtdt|jtdt|jtdt|jtdt|jtdt|jtdt|j tdt|j |jd vsJ|jd vsJ|j}t||t|j|ksJt|j |ksJt|j |ksJdS) Nrrrrrr r r) @)rrr) r<rrrrrrrr r rrr=r>s rr@zNVMMASharedLayout.__post_init__?s 02FtG^2_2_``` .0DTEZ0[0[\\\ F$8$C$CDDD L*>t*O*OPPP L*>t*O*OPPP N,@AR,S,STTT O-A$BT-U-UVVV K)=dn)M)MNNN$7777&*:::::yD$'''4$%%----4%&&$....4>""d******rc ~||j|j|j|j|j|j|jSr)get_nvmma_shared_layoutrrrrr r rr/s rr1zNVMMASharedLayout._to_irQsA..  #  ! O O     N   rc |rdnd}t||}t|} |r|dd|ddz}|d|z} | |jzdz} | dkr | dzdkrd} n&| dkr | dzdkrd} n| d kr | d zdkrd } nd} d} |ddD]}| |z} t|dks| dkrd} t| |j| ||||| S) zReturns an NVMMASharedLayout with default swizzling for a given shape. This picks the largest swizzle pattern compatible with the shape, which allows emitting the fewest TMA or MMA messages. r Nrrrrr)rrrrrr r r)rr=primitive_bitwidthr) block_shapedtyperrr r rpacking_factorrrcontig_dim_sizecontig_dim_bytesrflatten_outer_dimsizes rget_default_forz!NVMMASharedLayout.get_default_for\si)/a*; FF ;  B)!""- bqb0AAM'+n<*U-EEJ s " "'7#'='B'B!$    # #(82(=(B(B!#    # #(82(=(B(B!#  !" !#2#& & &D  %   {  a  #4q#8#8!"  1"5!!%'    rrCc Hd|jd|jd|jd|jd S)NNVMMA_rH_NVMMA)rrrrrs rr4zNVMMASharedLayout.mangles7s/ss$2Gss$/ss\`\kssssrc t|j|j|j|j|j|jrt|jnd|jrt|jnd|j rt|j ndfSr) rSrrrrrr rTr rrs rrUzNVMMASharedLayout.__hash__sT,d.CTYPTP_aeap151BLU4,---262DNU4-...$.2nFU4>***$HII Ir)FFNNN)r#r$r%r&rWrXrboolrr rrr rr@r1 staticmethodr rrLr4rUrYrZs@rrr'sI   IIIJJ(,L(49%,,,)-M8DI&---%)IxS ")))+++++$    qu"&& & & \& PtttttIIIIIIIrrceZdZUdZeed<eed<eed<eeed<dZeeeed<dZ eeeed<dZ eeeed <fd Z d Z d e fd ZdZxZS)SwizzledSharedLayouta Represents a generic swizzled shared memory layout. Args: vec (int): Vector width for swizzling. per_phase (int): Elements per swizzle phase. max_phase (int): Maximum number of swizzle phases. order (List[int]): Dimension ordering for swizzling. ctas_per_cga (Optional[List[int]]): CTAs per CGA grouping. cta_split_num (Optional[List[int]]): Split factors for CTAs. cta_order (Optional[List[int]]): CTA ordering. vec per_phase max_phaser:Nr r rctdt|jtdt|jtdt|jtdt|jtdt|jtdt|jtdt|j t|j}t||t|j|ksJt|j|ksJt|j |ksJdS)Nrrrr:r r r) r<rrrrrr:r r rr=rr>s rr@z"SwizzledSharedLayout.__post_init__s~ E#7#A#ABBB K)=dn)M)MNNN K)=dn)M)MNNN G%9$*%E%EFFF N,@AR,S,STTT O-A$BT-U-UVVV K)=dn)M)MNNN4:D$'''4$%%----4%&&$....4>""d******rc ~||j|j|j|j|j|j|jSr)get_swizzled_shared_layoutrrrr:r r rr/s rr1zSwizzledSharedLayout._to_irs?11 H N N J     N   rrCcd}d|jd|jd|jd||jd||jd||jd||jdS)NcZ|dSdtt|SrFrIrMs rrOz.SwizzledSharedLayout.mangle..stringifyrPrSSS_rH_SSS)rrrr:r r rrrOs rr4zSwizzledSharedLayout.mangles ) ) ) zdhzzzz$.zz99TZCXCXzz[d[deiev[w[wzz{D{DEIEW{X{Xzz[d[deies[t[tzzz zrc t|j|j|jt |j|jrt |jnd|jrt |jnd|jrt |jndfSr) rSrrrrTr:r r rrs rrUzSwizzledSharedLayout.__hash__sTXt~t~4:&&DDU(_d.?(@(@(@[_262DNU4-...$.2nFU4>***$HII Ir)r#r$r%r&rWrXrr rr rr@r1rLr4rUrYrZs@rrrs   HHHNNNNNN 9(,L(49%,,,)-M8DI&---%)IxS "))) + + + + +    zzzzzIIIIIIIrrceZdZUdZeeeed<eeed<dZeeeed<dZ eeeed<dZ eeeed<fdZ d Z d e fd Zd Zd ZxZS)PaddedSharedLayouta Represents a layout for the access to shared memory. Compared to SwizzledSharedLayout, it uses padding to avoid shared memory bank conflicts. After every interval tensor elements, the corresponding number of padding elements are inserted. If a position corresponds to multiple intervals, the padding amounts are summed. In the following example of a tensor, `eM` represents original elements in the and `pN` represents padded element. Before padding, the shared memory looks like: [e0, e1, e2, e3, e4, e5, e6, e7, ...] After padding with interval-padding list [[2, 1], [4, 2]], the shared memory will be [e0, e1, p0, e2, e3, p1, p2, p3, e4, e5, p4, e6, e7, p5, p6, p7, ...] Args: interval_padding_pairs (List[int]): List of [interval, padding] pair and both interval and padding must be powers of 2. order (List[int]): Order of logical tensor dimensions; fastest-varying first. ctas_per_cga (Optional[List[int]]): CTAs per CGA grouping. cta_split_num (Optional[List[int]]): Split factors for CTAs. cta_order (Optional[List[int]]): CTA ordering. interval_padding_pairsr:Nr r rc8tdt|jtdt |jtdt |jtdt |jtdt |j| dS)Nrr:r r r) r<rrrrr:r r rverifyr`s rr@z PaddedSharedLayout.__post_init__s 4mDD_6`6`aaa G%9$*%E%EFFF N,@AR,S,STTT O-A$BT-U-UVVV K)=dn)M)MNNN rct|j\}}||||j|j|j|jSr)ziprget_padded_shared_layoutr:r r r)rr0 intervalspaddingss rr1zPaddedSharedLayout._to_irsH!4#>? 8// 8TZQUQbdhdv04@@ @rrCc d}d||jd||jd||jd||jd||jd S)NcZ|dSdtt|SrFrIrMs rrOz,PaddedSharedLayout.mangle..stringifyrPr PaddedShared_rH _PaddedShared)rr:r r rrs rr4zPaddedSharedLayout.mangles ) ) ) Hyy)DEEHH RVR\H]H]HH`i`ijnj{`|`|HH@I@IJNJ\@]@]HH`i`ijnjx`y`yHHH Hrc|j}t|dks Jdtd|DsJt|\}}t t |}t|t|ksJdtfd|Ds Jdtfd|Ds Jdt|j}|dks Jd t||t|j|ksJt|j |ksJt|j |ksJdS) NrzVPaddedSharedLayout interval_padding_pairs must have at least one interval-padding pairc3<K|]}t|dkVdS)rN)r=).0pairs r z,PaddedSharedLayout.verify..s,44d3t99>444444rc&|dko ||dz zdkS)Nrr r()ns rz+PaddedSharedLayout.verify..s!a%".s-77==##777777rz;PaddedSharedLayout interval values must all be power of twoc3.K|]}|VdSrr(rs rrz,PaddedSharedLayout.verify..s-66==##666666rz:PaddedSharedLayout padding values must all be power of twoz*PaddedSharedLayout order must not be empty) rr=allrrsetr:rr r r)rpairsrrunique_intervalsrrs @rrzPaddedSharedLayout.verifys+5zzA~~~w~~~44e44444444!5k 8I//#$$I6666<< 7777Y77777vv9vvvv6666X66666tt8tttt4:axxxExxxD$'''4$%%----4%&&$....4>""d******rc :tttt|jt|j|jrt|jnd|jrt|jnd|jrt|jndfSr)rSrTrKrr:r r rrs rrUzPaddedSharedLayout.__hash__#sU3ud&ABBCC4:&&DDU(_d.?(@(@(@[_262DNU4-...$.2nFU4>***$HII Ir)r#r$r%r&rrWrXr rr rr@r1rLr4rrUrYrZs@rrrs >!cO+++ 9(,L(49%,,,)-M8DI&---%)IxS ")))@@@ HHHHH+++*IIIIIIIrrcdg|z}|s|Sd}|D]T}tdt|Dd}||}||xxdzcc<>|s|J||xxdzcc<U|S)Nr c3,K|]\}}|dk |VdS)rNr()rivs rrz bases_per_dim..5s*==$!Qa1ffAffff==rr)next enumerate)basesrskip_broadcastresult non_zero_idxroidxs r bases_per_dimr+sS4ZF  L & &==)E"2"2===tDD ?L 3KKK1 KKKK &+++ < A % Mrct|tr"t|jt |St|t t frt|j|S|jSr) isinstancerhrrkr=r\rvr9r^)rrms rr9r9Asa&122$V.E ;;; F[*:; < <$V]E222##rN)T) dataclassesrtypingrrtriton.language.corerrrtriton.runtime.jitr rrr,r6r\rhrvrrrrrrrr9r(rrrs!!!!!!!!!!!!!!TTTTTTTTTT111111777$$$$$$$$ $" $K K K K K %K K K \ $-----#---8 $2 2 2 2 2 /2 2 2 j $EEEEE(EEE8 $4   ,I,I,I,I,I.,I,I! ,I^$$$$$$$$ $eIeIeIeIeI eIeIeIP $4   <I<I<I<I<I<<I<I! <I~ $4   VIVIVIVIVIVIVI! VIt,$$$$$r