9iXSSKJr SSKJrJr SSKJrJrJr SSK J r Sr "SS5r \"SS 9"S S \ 55r \"SS 9"S S \ 55r\"SS 9"SS\ 55r\"SS 9"SS\ 55r\"SS 9"SS\ 55r\"SSS9"SS\ 55r"SS5r\ S5r\"SS 9"SS\55r\"SSS9"SS\55r\"SSS9"SS\55rS#S jrS!rg")$) dataclass)ListOptional)_unwrap_if_constexpr _unwrap_shapeconstexpr_type)constexpr_functioncRUR=(d S/U-nUR=(d S/U-nUR=(d [[ [ U555n[ RUSU5 [ RUSU5 [ RUSU5 g)N ctas_per_cga cta_split_num cta_order)r r rlistreversedrangeobject __setattr__)layoutrankr r rs k/var/www/html/land-doc-ocr/venv/lib/python3.13/site-packages/triton/experimental/gluon/language/_layouts.py_realize_cta_layoutrs&&41#*L((6QC$JM  ?D%+)>$?I v~|< v > v{I6c(\rSrSrSr\S5rSrg)DistributedLayoutz8 Base class for distributed memory layouts in Gluon IR. c[U5$Nrselfs rtypeDistributedLayout.type d##rN__name__ __module__ __qualname____firstlineno____doc__propertyr!__static_attributes__r$rrrr$$rrT)frozenc \rSrSrSrSrSrg) AutoLayoutc"UR5$r)get_auto_layoutr builders r_to_irAutoLayout._to_irs&&((rcg)NALr$rs rmangleAutoLayout.mangle srr$N)r&r'r(r)r6r:r,r$rrr0r0s )rr0c^\rSrSr%Sr\\\S'\\\S'\\\S'\\\S'Sr\ \\\S'Sr \ \\\S 'Sr \ \\\S 'U4S jr S r S \4SjrSrSrU=r$) BlockedLayout$a8 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 rct>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [TU]S[UR 55 [TU]S[UR55 [TU]S[UR55 [TU]S[UR55 [UR5n[X5 [UR5U:Xde[UR 5U:Xde[UR 5U:Xde[UR5U:Xde[UR5U:Xde[UR5U:Xdeg)Nr?r@rArBr r r) superrrr?r@rArBr r rlenrr r __class__s r __post_init__BlockedLayout.__post_init__:sg -/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 URURURURURUR UR UR5$r)get_blocked_layoutr?r@rArBr r rr4s rr6BlockedLayout._to_irLsP))   ! !    JJ       NN  rreturnc4SnU"UR5nU"UR5nU"UR5nU"UR5nU"UR5nU"UR 5nU"UR 5nSUSUSUSUSUSUSUS3$)NcHUcgSR[[U55$N_joinmapstrxs r stringify'BlockedLayout.mangle..stringifyYy88CQK( (rB)r?r@rArBr r r) r rYr?r@rArBr r rs rr:BlockedLayout.mangleWs ) $D$8$89$T%:%:;!$"4"45 $**% !2!23 !$"4"45 dnn- ?#1%5$6a awaP\~]^_l^mmnoxnyyz{{rc [[UR5[UR5[UR5[UR 5UR (a[UR 5OSUR(a[UR5OSUR(a[UR545$S45$r) hashtupler?r@rArBr r rrs r__hash__BlockedLayout.__hash__gs $&& ' $'' ( $$$ % $** (,(9(9E$## $t)-););E$$$ %%)^^E$.. !   :>   rr$r&r'r(r)r*rint__annotations__r rr rrHr6rVr:rar, __classcell__rGs@rr=r=$s #Y3i9 9(,L(49%,)-M8DI&-%)IxS ")+$  ||   rr=c\^\rSrSr%Sr\\S'\\S'U4SjrSr S\ 4Sjr S r S r U=r$) SliceLayoutsz 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. dimparentc>[TU]S[UR55 [TU]S[UR55 g)Nrkrl)rDrrrkrlr rGs rrHSliceLayout.__post_init__s5 E#7#AB H&:4;;&GHrclURURURRU55$r)get_slice_layoutrkrlr6r4s rr6SliceLayout._to_irs.'' HH KK  w '  rrMcXSURSURR5S3$)NSLrR)rkrlr:rs rr:SliceLayout.mangles)DHH:Qt{{1134B77rcD[URUR45$r)r_rkrlrs rraSliceLayout.__hash__sTXXt{{+,,rr$r&r'r(r)r*rdrerrHr6rVr:rar,rfrgs@rririss6 H I 88--rric^\rSrSr%Sr\\\\S'\\\\S'\\\\S'\\\\S'\\\S'U4SjrS r S r S r S r U=r $) 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_basesshapec>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [TU]S[UR 55 [TU]S[UR55 [UR5nURHn[U5U:XaMe URHn[U5U:XaMe UR Hn[U5U:XaMe UR Hn[U5U:XaMe g)Nr|r}r~rr) rDrrr|r}r~rrrE)r rbasisrGs rrH%DistributedLinearLayout.__post_init__s  Kt~~)FG L-*HI L-*HI M=9I9I+JK G]4::%>?4::^^Eu:% %%$__Eu:% %%%__Eu:% %%%%%Eu:% %%&rcURURURURURUR 5$r)get_distributed_linear_layoutr|r}r~rrr4s rr6DistributedLinearLayout._to_irs<44T^^T__VZVeVegkgwgw59ZZA Arc SURSURSURSURSURS3 $)NDLLrR)r|r}r~rrrs rr:DistributedLinearLayout.manglesIT^^$Adoo%6a7H$JZJZI[[\]a]g]g\hhkllrc P[[[[UR55[[[UR55[[[UR 55[[[UR 55[UR545$r)r_r`rUr|r}r~rrrs rra DistributedLinearLayout.__hash__sn #eT^^, - #eT__- . #eT__- . #eT--. / $**     rr$)r&r'r(r)r*rrdrerHr6r:rar,rfrgs@rrzrzsg DIT#YT#Yd3i 9&$Am  rrzcf^\rSrSr%Sr\\S'\\S'\\S'U4SjrSr S\ 4S jr S r S r U=r$) DotOperandLayoutz 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_indexrlk_widthc>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 g)Nrrlr)rDrrrrlrrns rrHDotOperandLayout.__post_init__sP O-A$BTBT-UV H&:4;;&GH I';DLL'IJrcURURURRU5UR5$r)get_dot_operand_layoutrrlr6rr4s rr6DotOperandLayout._to_irs4--d.@.@$++BTBTU\B]_c_k_kllrrMcrSURSURR5SURS3$)NDOrR)rrlr:rrs rr:DotOperandLayout.mangles6D&&'q););)=(>a ~RPPrcZ[URURUR45$r)r_rrlrrs rraDotOperandLayout.__hash__s"T''dllCDDrr$rxrgs@rrrsA  LK mQQEErr)r.eqc^\rSrSr%Sr\\\S'\\\S'\\\S'Sr\ \\\S'Sr \ \\\S'Sr \ \\\S 'U4S jr S r S \4S jrSrSrU=r$)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. versionrA instr_shapeNr r rc>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [TU]S[UR 55 [TU]S[UR55 [TU]S[UR55 [UR5n[X5 [UR 5U:Xde[UR5U:Xde[UR5U:Xdeg)NrrArr r r) rDrrrrArr r rrErrFs rrH$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***rcURURURURURUR UR 5$r)get_mma_layoutrrAr r rrr4s rr6NVMMADistributedLayout._to_irsE%%dllD4F4FHYHY[_[m[m&*nnd6F6FH HrrMc SURSURSURSURSURSUR S3 $)NMMA_rR_MMA)rrArr r rrs rr:NVMMADistributedLayout.mangle sqdll^1T%7%7$8$:J:J9K1TM^M^L__`aeasas`ttuvzwEwEvFFJK Krc [[UR5[UR5[UR5UR (a[UR 5OSUR (a[UR 5OSUR(a[UR545$S45$r)r_r`rrArr r rrs rraNVMMADistributedLayout.__hash__ sU4<<(%0B0B*C4++,$J[J[eD4E4E.Fae262D2DU4--.$.2nnU4>>*HI ICGHI Irr$rcrgs@rrrs #Y9c(,L(49%,)-M8DI&-%)IxS ") +HKKIIrrc(\rSrSrSr\S5rSrg) SharedLayoutiz3 Base class for shared memory layouts in Gluon IR. c[U5$rrrs rr!SharedLayout.typer#rr$Nr%r$rrrrr-rrcUnUbC[U5[U5:Xde[[U55HnX#==X-ss'M U$r)rEr)rr shape_per_ctarks r_get_shape_per_ctarsOM =!SZ///]+,C  -"4 4 - rc^\rSrSr%Sr\\S'\\S'\\S'Sr\\S'Sr \\S'S r \ \ \\S 'S r \ \ \\S 'S r\ \ \\S 'U4S jrSr\\SSj55rS\4SjrSrSrU=r$)NVMMASharedLayouti'a 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 rcN>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [TU]S[UR 55 [TU]S[UR55 [TU]S[UR55 [TU]S[UR55 [TU]S[UR55 URS ;deURS ;deUR n[X5 [UR5U:Xde[UR5U:Xde[UR5U:Xdeg) Nrrrrrr r r)r @)rrr) rDrrrrrrrr r rrrErFs rrHNVMMASharedLayout.__post_init__?sX 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 URURURURURUR UR UR5$r)get_nvmma_shared_layoutrrrrr r rr4s rr6NVMMASharedLayout._to_irQsN..  # #  ! ! OO OO       NN  rc U(aSOSn[X5n[U5n U(a USSUSS-nUSU-n XR-S-n U S:a U S-S:XaSn O&U S:a U S-S:XaSn OU S :a U S -S:XaS n OSn Sn USSHnX-n M [U5S:dU S:aSn [U URU UUUUUS 9$) 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)rrEprimitive_bitwidthr) block_shapedtyperrr r rpacking_factorrrcontig_dim_sizecontig_dim_bytesrflatten_outer_dimsizes rget_default_for!NVMMASharedLayout.get_default_for\s)a*;F ; )!"- bq0AAM'+n<*-E-EEJ s "'7#'='B!$   #(82(=(B!#   #(82(=(B!# !" !#2&D  % ' { a #4q#8!"  1"55!!%'  rrMc pSURSURSURSURS3 $)NNVMMA_rR_NVMMA)rrrrrs rr:NVMMASharedLayout.mangles@//0$2G2G1H$//IZZ[\`\k\k[llrssrc z[URURURURUR UR (a[UR 5OSUR(a[UR5OSUR(a[UR545$S45$r) r_rrrrrr r`r rrs rraNVMMASharedLayout.__hash__sT,,d.C.CTYYPTP_P_aeapap151B1BU4,,-262D2DU4--.$.2nnU4>>*HI ICGHI Irr$)FFNNN)r&r'r(r)r*rdrerboolrr rrr rrHr6 staticmethodr rrVr:rar,rfrgs@rrr's  IJJ(,L(49%,)-M8DI&-%)IxS ")+$  qu"&& & PttIIrrc^\rSrSr%Sr\\S'\\S'\\S'\\\S'Sr\ \\\S'Sr \ \\\S 'Sr \ \\\S 'U4S jr S r S \4SjrSrSrU=r$)SwizzledSharedLayoutia 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_phaserBNr r rc>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [TU]S[UR 55 [TU]S[UR55 [TU]S[UR55 [TU]S[UR55 [UR 5n[X5 [UR5U:Xde[UR5U:Xde[UR5U:Xdeg)NrrrrBr r r) rDrrrrrrBr r rrErrFs rrH"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 URURURURURUR UR UR5$r)get_swizzled_shared_layoutrrrrBr r rr4s rr6SwizzledSharedLayout._to_irsJ11 HH NN NN JJ       NN  rrMcSnSURSURSURSU"UR5SU"UR5SU"UR 5SU"UR 5S3$)NcHUcgSR[[U55$rPrSrWs rrY.SwizzledSharedLayout.mangle..stringifyr[rSSS_rR_SSS)rrrrBr r rr rYs rr:SwizzledSharedLayout.mangles ) dhhZq 0$..1A9TZZCXBYYZ[deievev[wZxxy{DEIEWEW{XzYYZ[deieses[tZuuyz zrc v[URURUR[ UR 5UR (a[ UR 5OSUR(a[ UR5OSUR(a[ UR545$S45$r) r_rrrr`rBr r rrs rraSwizzledSharedLayout.__hash__sTXXt~~t~~4::&DDUDUd.?.?(@[_262D2DU4--.$.2nnU4>>*HI ICGHI Irr$)r&r'r(r)r*rdrerr rr rrHr6rVr:rar,rfrgs@rrrs  HNN 9(,L(49%,)-M8DI&-%)IxS ") +  zzIIrrc^\rSrSr%Sr\\\\S'\\\S'Sr\ \\\S'Sr \ \\\S'Sr \ \\\S'U4S jr S r S \4S jrS rSrSrU=r$)PaddedSharedLayoutia6 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_pairsrBNr r rc>[TU]S[UR55 [TU]S[ UR 55 [TU]S[ UR 55 [TU]S[ UR55 [TU]S[ UR55 UR5 g)NrrBr r r) rDrrrrrBr r rverifyrns rrH PaddedSharedLayout.__post_init__s 4mDD_D_6`a G%9$**%EF N,@ARAR,ST O-A$BTBT-UV K)=dnn)MN rc[UR6up#URX#URURUR UR 5$r)ziprget_padded_shared_layoutrBr r r)r r5 intervalspaddingss rr6PaddedSharedLayout._to_irsK!4#>#>? // TZZQUQbQbdhdvdv04@ @rrMc SnSU"UR5SU"UR5SU"UR5SU"UR5SU"UR5S3 $)NcHUcgSR[[U55$rPrSrWs rrY,PaddedSharedLayout.mangle..stringifyr[r PaddedShared_rR _PaddedShared)rrBr r rrs rr:PaddedSharedLayout.mangles ) y)D)DEFa RVR\R\H]G^^_`ijnj{j{`|_}}~@IJNJ\J\@]^^_`ijnjxjx`y_zzGH Hrcr^URn[U5S:dS5e[SU55(de[U6up#[ [ U55n[U5[U5:XdeSm[U4SjU55(dS5e[U4SjU55(dS5e[UR 5nUS:dS 5e[X5 [UR5U:Xde[UR5U:Xde[UR5U:Xdeg) NrzVPaddedSharedLayout interval_padding_pairs must have at least one interval-padding pairc3># UHn[U5S:Hv M g7f)rN)rE).0pairs r ,PaddedSharedLayout.verify..s4ed3t9>esc.US:=(a XS- -S:H$)Nrr r$)ns r+PaddedSharedLayout.verify..s!a%"# UH nT"U5v M g7frr$rr is_power_of_2s rrrs7Y=##Yz;PaddedSharedLayout interval values must all be power of twoc34># UH nT"U5v M g7frr$rs rrrs6X=##Xr z:PaddedSharedLayout padding values must all be power of twoz*PaddedSharedLayout order must not be empty) rrEallrrsetrBrr r r)r pairsrrunique_intervalsrrs @rrPaddedSharedLayout.verifys++5zA~www~4e44444!5k I/#$I666< 7Y777v9vv76X666t8tt64::axEEExD'4$$%---4%%&$...4>>"d***rc x[[[[UR55[UR5UR (a[UR 5OSUR (a[UR 5OSUR(a[UR545$S45$r)r_r`rUrrBr r rrs rraPaddedSharedLayout.__hash__#sU3ud&A&ABC4::&DDUDUd.?.?(@[_262D2DU4--.$.2nnU4>>*HI ICGHI Irr$)r&r'r(r)r*rrdrer rr rrHr6rVr:rrar,rfrgs@rrrs>!cO+ 9(,L(49%,)-M8DI&-%)IxS ")@ HH+*IIrrcS/U-nU(dU$SnUHLn[S[U55S5nUbUnX6==S-ss'M2U(aM;UceX4==S-ss'MN U$)Nr c3:# UHupUS:wdM Uv M g7f)rNr$)rivs rr bases_per_dim..5s="2$!a1fAA"2s  r)next enumerate)basesrskip_broadcastresult non_zero_idxridxs r bases_per_dimr+swS4ZF  L=)E"2=tD ?L K1 K+ ++  A %  Mrc[U[5(a[UR[ U55$[U[ [ 45(a[URU5$UR$r) isinstancerzrr~rErirrArl)rrs rrArAAsW&122V..E ;; F[*:; < <V]]E22###rN)T) dataclassesrtypingrrtriton.language.corerrrtriton.runtime.jitr rrr0r=rirzrrrrrrrrrAr$rrr%s!!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