c iX0ddlmZddlmZmZddlmZmZmZddl 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!Zy")$) dataclass)ListOptional)_unwrap_if_constexpr _unwrap_shapeconstexpr_type)constexpr_functionc4|jxsdg|z}|jxsdg|z}|jxstt t |}t j|d|t j|d|t j|d|y)N ctas_per_cga cta_split_num cta_order)r r rlistreversedrangeobject __setattr__)layoutrankr r rs e/var/www/html/engine/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/_layouts.py_realize_cta_layoutrs&&41#*L((6QC$JM  ?D%+)>$?I v~|< v > v{I6c eZdZdZedZy)DistributedLayoutz@ Base class for distributed memory layouts in Gluon IR. ct|SNrselfs rtypezDistributedLayout.type d##rN__name__ __module__ __qualname____doc__propertyr rrrr$$rrT)frozenceZdZdZdZy) AutoLayoutc"|jSr)get_auto_layoutrbuilders r_to_irzAutoLayout._to_irs&&((rcy)NALr(rs rmanglezAutoLayout.mangle srN)r#r$r%r1r4r(rrr,r,s )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 rcjt|dt|jt|dt|jt|dt|j t|dt|j t|dt|jt|dt|jt|dt|jt|j}t||t|j|k(sJt|j |k(sJt|j |k(sJt|j|k(sJt|j|k(sJt|j|k(sJy)Nr7r8r9r:r r r) superrrr7r8r9r:r r rlenrrr __class__s r __post_init__zBlockedLayout.__post_init__:si -/CDDXDX/YZ .0DTEZEZ0[\ O-A$BTBT-UV G%9$**%EF N,@ARAR,ST O-A$BTBT-UV K)=dnn)MN4''(D$'4(()T1114%%&$...4::$&&&4$$%---4%%&$...4>>"d***rc |j|j|j|j|j|j |j |jSr)get_blocked_layoutr7r8r9r:r r rr/s rr1zBlockedLayout._to_irLsP))   ! !    JJ       NN  rreturnc4d}||j}||j}||j}||j}||j}||j }||j }d|d|d|d|d|d|d|dS)NcF|ydjtt|SN_joinmapstrxs r stringifyz'BlockedLayout.mangle..stringifyYy88CQK( (rB)r7r8r9r:r r r) rrOr7r8r9r:r r rs rr4zBlockedLayout.mangleWs ) $D$8$89$T%:%:;!$"4"45 $**% !2!23 !$"4"45 dnn- ?#1%5$6a awaP\~]^_l^mmnoxnyyz{{rc tt|jt|jt|jt|j |j rt|j nd|jrt|jnd|jrt|jfSdfSr) hashtupler7r8r9r:r r rrs r__hash__zBlockedLayout.__hash__gs $&& ' $'' ( $$$ % $** (,(9(9E$## $t)-););E$$$ %%)^^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. dimparentct|dt|jt|dt|jy)Nr]r^)r<rrr]r^rr?s rr@zSliceLayout.__post_init__s5 E#7#AB H&:4;;&GHrcl|j|j|jj|Sr)get_slice_layoutr]r^r1r/s rr1zSliceLayout._to_irs.'' HH KK  w '  rrCcXd|jd|jjdS)NSLrH)r]r^r4rs rr4zSliceLayout.mangles)DHH:Qt{{1134B77rcDt|j|jfSr)rSr]r^rs rrUzSliceLayout.__hash__sTXXt{{+,,r r#r$r%r&rWrXrr@r1rLr4rUrYrZs@rr\r\ss1 H I 88-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_basesshapect|dt|jt|dt|jt|dt|j t|dt|j t|dt|jt|j}|jD]}t||k(rJ|jD]}t||k(rJ|j D]}t||k(rJ|j D]}t||k(rJy)Nrirjrkrlrm) r<rrrirjrkrlrmr=)rrbasisr?s rr@z%DistributedLinearLayout.__post_init__s  Kt~~)FG L-*HI L-*HI M=9I9I+JK G]4::%>?4::^^ &Eu:% %% &__ &Eu:% %% &__ &Eu:% %% &%% &Eu:% %% &rc|j|j|j|j|j|j Sr)get_distributed_linear_layoutrirjrkrlrmr/s rr1zDistributedLinearLayout._to_irs<44T^^T__VZVeVegkgwgw59ZZA Arc d|jd|jd|jd|jd|jd S)NDLLrH)rirjrkrlrmrs rr4zDistributedLinearLayout.manglesIT^^$Adoo%6a7H$JZJZI[[\]a]g]g\hhkllrc Ptttt|jttt|jttt|j ttt|j t|jfSr)rSrTrKrirjrkrlrmrs rrUz DistributedLinearLayout.__hash__sn #eT^^, - #eT__- . #eT__- . #eT--. / $**     r) r#r$r%r&rrWrXr@r1r4rUrYrZs@rrhrhsb DIT#YT#Yd3i 9&$Am 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_widthct|dt|jt|dt|jt|dt|j y)Nrwr^rx)r<rrrwr^rxr`s rr@zDotOperandLayout.__post_init__sP O-A$BTBT-UV H&:4;;&GH I';DLL'IJrc|j|j|jj||jSr)get_dot_operand_layoutrwr^r1rxr/s rr1zDotOperandLayout._to_irs4--d.@.@$++BTBTU\B]_c_k_kllrrCcrd|jd|jjd|jdS)NDOrH)rwr^r4rxrs rr4zDotOperandLayout.mangles6D&&'q););)=(>a ~RPPrcZt|j|j|jfSr)rSrwr^rxrs rrUzDotOperandLayout.__hash__s"T''dllCDDrrfrZs@rrvrvs;  LK mQQErrv)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 rct|dt|jt|dt|jt|dt|j t|dt|j t|dt|jt|dt|jt|j}t||t|j |k(sJt|j|k(sJt|j|k(sJy)Nrr9rr r r) r<rrrr9rr r rr=rr>s rr@z$NVMMADistributedLayout.__post_init__s I';DLL'IJ O-A$BTBT-UV M+?@P@P+QR N,@ARAR,ST O-A$BTBT-UV K)=dnn)MN4%%&D$'4$$%---4%%&$...4>>"d***rc|j|j|j|j|j|j |j Sr)get_mma_layoutrr9r r rrr/s rr1zNVMMADistributedLayout._to_irsE%%dllD4F4FHYHY[_[m[m&*nnd6F6FH HrrCc d|jd|jd|jd|jd|jd|j d S)NMMA_rH_MMA)rr9rr r rrs rr4zNVMMADistributedLayout.mangle sqdll^1T%7%7$8$:J:J9K1TM^M^L__`aeasas`ttuvzwEwEvFFJK Krc ftt|jt|jt|j|j rt|j nd|j rt|j nd|jrt|jfSdfSr)rSrTrr9rr r rrs rrUzNVMMADistributedLayout.__hash__ sU4<<(%0B0B*C4++,$J[J[eD4E4E.Fae262D2DU4--.$.2nnU4>>*HI ICGHI IrrVrZs@rrrs #Y9c(,L(49%,)-M8DI&-%)IxS ") +HKKIrrc eZdZdZedZy) SharedLayoutz; Base class for shared memory layouts in Gluon IR. ct|Srrrs rr zSharedLayout.typer!rNr"r(rrrrr)rrc|}|Bt|t|k(sJtt|D]}||xx||zcc<|Sr)r=r)rmr shape_per_ctar]s r_get_shape_per_ctarsXM =!SZ///]+, 5C # -"4 4  5 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 rcBt|dt|jt|dt|jt|dt|j t|dt|j t|dt|jt|dt|jt|dt|jt|dt|j|jd vsJ|jd vsJ|j }t||t|j|k(sJt|j|k(sJt|j|k(sJy) Nrrrrrr r r) @)rrr) r<rrrrrrrr r rrr=r>s rr@zNVMMASharedLayout.__post_init__?sZ 02FtG^G^2_` .0DTEZEZ0[\ F$8$CD L*>t*OP L*>t*OP N,@ARAR,ST O-A$BTBT-UV K)=dnn)MN$$777&&*::::yyD$'4$$%---4%%&$...4>>"d***rc |j|j|j|j|j|j |j |jSr)get_nvmma_shared_layoutrrrrr r rr/s rr1zNVMMASharedLayout._to_irQsN..  # #  ! ! OO OO       NN  rc t|rdnd}t||}t|} |r |dd|ddz}|d|z} | |jzdz} | dk\r | dzdk(rd} n"| dk\r | dzdk(rd} n| d k\r | d zdk(rd } 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\s')a*; F ; )!"- bq0AAM'+n<*U-E-EEJ s "'7#'='B!$   #(82(=(B!#   #(82(=(B!# !" !#2& &D  %  & { a #4q#8!"  1"55!!%'  rrCc pd|jd|jd|jd|jd S)NNVMMA_rH_NVMMA)rrrrrs rr4zNVMMASharedLayout.mangles@//0$2G2G1H$//IZZ[\`\k\k[llrssrc \t|j|j|j|j|j |j rt|j nd|jrt|jnd|jrt|jfSdfSr) rSrrrrrr rTr rrs rrUzNVMMASharedLayout.__hash__sT,,d.C.CTYYPTP_P_aeapap151B1BU4,,-262D2DU4--.$.2nnU4>>*HI ICGHI Ir)FFNNN)r#r$r%r&rWrXrboolrr rrr rr@r1 staticmethodr rrLr4rUrYrZs@rrr's  IJJ(,L(49%,)-M8DI&-%)IxS ")+$  qu"&& & PttIrrceZdZUdZeed<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 rct|dt|jt|dt|jt|dt|j t|dt|j t|dt|jt|dt|jt|dt|jt|j }t||t|j|k(sJt|j|k(sJt|j|k(sJy)Nrrrr:r r r) r<rrrrrr:r r rr=rr>s rr@z"SwizzledSharedLayout.__post_init__s E#7#AB K)=dnn)MN K)=dnn)MN G%9$**%EF N,@ARAR,ST O-A$BTBT-UV K)=dnn)MN4::D$'4$$%---4%%&$...4>>"d***rc |j|j|j|j|j|j |j |jSr)get_swizzled_shared_layoutrrrr:r r rr/s rr1zSwizzledSharedLayout._to_irsJ11 HH NN NN JJ       NN  rrCcd}d|jd|jd|jd||jd||jd||j d||j dS)NcF|ydjtt|SrFrIrMs rrOz.SwizzledSharedLayout.mangle..stringifyrPrSSS_rH_SSS)rrrr:r r rrrOs rr4zSwizzledSharedLayout.mangles ) dhhZq 0$..1A9TZZCXBYYZ[deievev[wZxxy{DEIEWEW{XzYYZ[deieses[tZuuyz zrc Xt|j|j|jt |j |j rt |j nd|jrt |jnd|jrt |jfSdfSr) rSrrrrTr:r r rrs rrUzSwizzledSharedLayout.__hash__sTXXt~~t~~4::&DDUDUd.?.?(@[_262D2DU4--.$.2nnU4>>*HI ICGHI Ir)r#r$r%r&rWrXrr rr rr@r1rLr4rUrYrZs@rrrsz  HNN 9(,L(49%,)-M8DI&-%)IxS ") +  zzIrrceZdZUdZeeeed<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 rct|dt|jt|dt |j t|dt |j t|dt |jt|dt |j|jy)Nrr:r r r) r<rrrrr:r r rverifyr`s rr@z PaddedSharedLayout.__post_init__s 4mDD_D_6`a G%9$**%EF N,@ARAR,ST O-A$BTBT-UV K)=dnn)MN rct|j\}}|j|||j|j|j |j Sr)ziprget_padded_shared_layoutr:r r r)rr0 intervalspaddingss rr1zPaddedSharedLayout._to_irsO!4#>#>? 8// 8TZZQUQbQbdhdvdv04@ @rrCc d}d||jd||jd||jd||jd||jd S)NcF|ydjtt|SrFrIrMs rrOz,PaddedSharedLayout.mangle..stringifyrPr PaddedShared_rH _PaddedShared)rr:r r rrs rr4zPaddedSharedLayout.mangles ) y)D)DEFa RVR\R\H]G^^_`ijnj{j{`|_}}~@IJNJ\J\@]^^_`ijnjxjx`y_zzGH HrcH|j}t|dkDsJdtd|DsJt|\}}t t |}t|t|k(sJdtfd|DsJdtfd|DsJdt|j }|dkDsJd t||t|j|k(sJt|j|k(sJt|j|k(sJy) NrzVPaddedSharedLayout interval_padding_pairs must have at least one interval-padding pairc38K|]}t|dk(yw)rN)r=).0pairs r z,PaddedSharedLayout.verify..s4d3t9>4sc&|dkDxr ||dz zdk(S)Nrr r()ns rz+PaddedSharedLayout.verify..s!a%".s7=#7z;PaddedSharedLayout interval values must all be power of twoc3.K|] }|ywrr(rs rrz,PaddedSharedLayout.verify..s6=#6rz:PaddedSharedLayout padding values must all be power of twoz*PaddedSharedLayout order must not be empty) rr=allrrsetr:rr r r)rpairsrrunique_intervalsrrs @rrzPaddedSharedLayout.verifys++5zA~www~4e4444!5k 8I/#$I666< 7Y77v9vv76X66t8tt64::axEEExD$'4$$%---4%%&$...4>>"d***rc Ztttt|jt|j|j rt|j nd|j rt|j nd|jrt|jfSdfSr)rSrTrKrr:r r rrs rrUzPaddedSharedLayout.__hash__#sU3ud&A&ABC4::&DDUDUd.?.?(@[_262D2DU4--.$.2nnU4>>*HI ICGHI Ir)r#r$r%r&rrWrXr rr rr@r1rLr4rrUrYrZs@rrrs>!cO+ 9(,L(49%,)-M8DI&-%)IxS ")@ HH+*Irrcdg|z}|s|Sd}|D]D}tdt|Dd}||}||xxdzcc<1|r4|J||xxdzcc<F|S)Nr c32K|]\}}|dk7s |yw)rNr()rivs rrz bases_per_dim..5s=$!Qa1fA=s r)next enumerate)basesrskip_broadcastresult non_zero_idxroidxs r bases_per_dimr+s~S4ZF  L &=)E"2=tD ?L 3K1 K+ ++ < A %  & Mrct|trt|jt |St|t t frt|j|S|jSr) isinstancerhrrkr=r\rvr9r^)rrms rr9r9AsQ&12V..E ;; F[*:; <V]]E22###rN)T) dataclassesrtypingrrtriton.language.corerrrtriton.runtime.jitr rrr,r6r\rhrvrrrrrrrr9r(rrrs!!TT17$$ $" $K %K K \ $-#--8 $2 /2 2 j $E(EE8 $4 ,I.,I!,I^$$ $eI eIeIP $4 <I<<I!<I~ $4 VIVI!VIt,$r