1jZeddlmZmZddlZddlZddlmZddlmZm Z m Z ddl m Z ddl mZGddZed Gd d eZed Gd deZed GddeZed GddeZed GddeZed GddeZed d GddeZGddZe dZed GddeZed d Gdd eZed d Gd!d"eZed Gd#d$eZd'd%Zd&Zy)() dataclassfieldN)List)_unwrap_if_constexpr _unwrap_shapeconstexpr_type)constexpr_function)gluon_irc\eZdZdZedZedZdeede fdZ deede fdZ y) DistributedLayoutz@ Base class for distributed memory layouts in Gluon IR. ct|SNrselfs j/var/www/html/ia-enhanced/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/_layouts.pytypezDistributedLayout.type d##ctd)Nz-DistributedLayout subclasses must define rank)NotImplementedErrorrs rrankzDistributedLayout.ranks!"QRRrshapereturnc htj||Dcgc] }t|c}dScc}wNFr get_layout_viewrrrss rformat_tensor_viewz$DistributedLayout.format_tensor_view-''PU.V1/CA/F.VX]^^.V/ c htj||Dcgc] }t|c}dScc}wNTrrs rformat_hardware_viewz&DistributedLayout.format_hardware_view-''PU.V1/CA/F.VX\]].Vr#N) __name__ __module__ __qualname____doc__propertyrrlistintstrr!r&rrr r se$$SS_S _c_^$s)^^rr T)frozenc(eZdZdZdZedZy) AutoLayoutc"|jSr)get_auto_layoutrbuilders r_to_irzAutoLayout._to_ir"s&&((rcy)NALr0rs rmanglezAutoLayout.mangle%rctd)NzAutoLayout has no rank ValueErrorrs rrzAutoLayout.rank(s122rNr(r)r*r8r;r,rr0rrr3r3s )33rr3c(eZdZdZdZedZy)CoalescedLayoutc"|jSr)get_coalesced_layoutr6s rr8zCoalescedLayout._to_ir0s++--rcy)NCLr0rs rr;zCoalescedLayout.mangle3r<rctd)NzCoalescedLayout has no rankr>rs rrzCoalescedLayout.rank6s677rNr@r0rrrBrB-s .88rrBceZdZUdZeeed<eeed<eeed<eeed<eeZ eeeed<fdZ d Z d e fd Z d Zed ZxZS) 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. cga_layout (Optional[List[List[int]]]): Bases describing how CTAs tile each dimension. size_per_threadthreads_per_warp warps_per_ctaorderdefault_factory cga_layoutc*t|dt|jt|dt|jt|dt|j t|dt|j t|dt|jt|j}t|j|k(sJt|j |k(sJt|j |k(sJy)NrJrKrLrMrP) super __setattr__rrJrKrLrMrPlen)rr __class__s r __post_init__zBlockedLayout.__post_init__Ms -/CDDXDX/YZ .0DTEZEZ0[\ O-A$BTBT-UV G%9$**%EF L*>t*OP4''(4(()T1114%%&$...4::$&&&rc|j|j|j|j|j|j Sr)get_blocked_layoutrJrKrLrMrPr6s rr8zBlockedLayout._to_irYs>))   ! !    JJ OO   rrc d}||j}||j}||j}||j}|jr"dj d|jDnd}d|d|d|d|d|d S)NcF|ydjtt|SN_joinmapr/xs r stringifyz'BlockedLayout.mangle..stringifydy88CQK( (rr]c3ZK|]#}djtt|%yw~Nr^.0vecs r z'BlockedLayout.mangle..mQ#chhs3}5Q)+r\B)rJrKrLrMrPr_)rrcrJrKrLrMrPs rr;zBlockedLayout.manglebs ) $D$8$89$T%:%:;!$"4"45 $**%UYUdUdSXXQQQjl ?#1%5$6a awaPZ|[\]]rc tt|jt|jt|jt|j td|j DfS)Nc32K|]}t|ywrtuplerhs rrkz)BlockedLayout.__hash__..rs-TSeCj-T)hashrrrJrKrLrMrPrs r__hash__zBlockedLayout.__hash__ps]U4//0%8M8M2NPUVZVhVhPi4::&-TDOO-T(TVW Wrc,t|jSr)rTrMrs rrzBlockedLayout.rankt4::rr(r)r*r+rr.__annotations__rr-rPrVr8r/r;rur,r __classcell__rUs@rrIrI;s{ #Y3i9 9"'"=JT#Y= '  ^ ^WrrIcpeZdZUdZeed<eed<fdZdZde fdZ dZ e d Z e d ZxZS) 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)rRrSrr~rrrUs rrVzSliceLayout.__post_init__s5 E#7#AB H&:4;;&GHrcl|j|j|jj|Sr)get_slice_layoutr~rr8r6s rr8zSliceLayout._to_irs.'' HH KK  w '  rrcXd|jd|jjdS)NSLr])r~rr;rs rr;zSliceLayout.mangles)DHH:Qt{{1134B77rcDt|j|jfSr)rtr~rrs rruzSliceLayout.__hash__sTXXt{{+,,rc4|jjdz S)Nrrrs rrzSliceLayout.ranks{{!##rc|jj}|sgS|jj}d|jcxkr|ksJJ|Dcgc]$}|d|j||jdzdz&c}Scc}wNrr)rrPrr~)rparent_cga_layoutrbasiss rrPzSliceLayout.cga_layoutsy KK22 I{{DHH#t#####EVWEitxx 5A#77WWWs)A;r(r)r*r+r.ryr rVr8r/r;rur,rrPrzr{s@rr}r}ys[ H I 88-$$XXrr}ceZdZUdZeeeed<eeeed<eeeed<eeeed<eeed<fdZdZd Z d Z e 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_basesrct|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)Nrrrrr) rRrSrrrrrrrTrrrrUs rrVz%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_layoutrrrrrr6s rr8zDistributedLinearLayout._to_irs<44T^^T__VZVeVegkgwgw59ZZA Arc d|jd|jd|jd|jd|jd S)NDLLr])rrrrrrs rr;zDistributedLinearLayout.manglesIT^^$Adoo%6a7H$JZJZI[[\]a]g]g\hhkllrc Ptttt|jttt|jttt|j ttt|j t|jfSr)rtrrr`rrrrrrs rruz DistributedLinearLayout.__hash__sn #eT^^, - #eT__- . #eT__- . #eT--. / $**     rc,t|jSr)rTrrs rrzDistributedLinearLayout.rankrwr)r(r)r*r+rr.ryrVr8r;rur,rrzr{s@rrrsv DIT#YT#Yd3i 9&$Am rrczeZdZUdZeed<eed<eed<fdZdZde fdZ d Z e d Z e d ZxZS) 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_indexrk_widthct|dt|jt|dt|jt|dt|j y)Nrrr)rRrSrrrrrs rrVzDotOperandLayout.__post_init__sP O-A$BTBT-UV H&:4;;&GH I';DLL'IJrc|j|j|jj||jSr)get_dot_operand_layoutrrr8rr6s rr8zDotOperandLayout._to_irs4--d.@.@$++BTBTU\B]_c_k_kllrrcrd|jd|jjd|jdS)NDOr])rrr;rrs rr;zDotOperandLayout.mangles6D&&'q););)=(>a ~RPPrcZt|j|j|jfSr)rtrrrrs rruzDotOperandLayout.__hash__s"T''dllCDDrc.|jjSrrrs rrzDotOperandLayout.ranks{{rcVtt|jdgxsg}|sgS|jjt fd|DsJ|j dk(rdz ndz }d|cxkrksJJg}|D]#}t |}d||<|j|%|S)NrPc3:K|]}t|k(ywrrT)rirrs rrkz.DotOperandLayout.cga_layout..sE%3u:%Esrr)rgetattrrrallrr-append)rrk_dimderivedr new_basisrs @rrPzDotOperandLayout.cga_layouts0lTV1WX^\^ I{{E3DEEEE ..!3qE D     & &EU I Ie  NN9 % &rrr{s@rrrsc  LK mQQE  rr)r1eqceZdZUdZeeed<eeed<eeed<eeZ eeeed<fdZ dZ d e fd Z d Zed 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. cga_layout (Optional[List[List[int]]]): Bases describing CTA tiling. versionrL instr_shaperNrPct|dt|jt|dt|jt|dt|j t|dt|j y)NrrLrrP)rRrSrrrLrrPrs rrVz$NVMMADistributedLayout.__post_init__"sk I';DLL'IJ O-A$BTBT-UV M+?@P@P+QR L*>t*OPrcz|j|j|j|j|jSr)get_mma_layoutrrLrPrr6s rr8zNVMMADistributedLayout._to_ir(s5%% LL    OO      rrc |jr"djd|jDnd}d|jd|jd|jd|d S)Nr]c3ZK|]#}djtt|%ywrfr^rhs rrkz0NVMMADistributedLayout.mangle..1rlrmr\MMA__MMA)rPr_rrLrrrPs rr;zNVMMADistributedLayout.mangle0sYUYUdUdSXXQQQjl dll^1T%7%7$8$:J:J9K1ZLX\]]rc tt|jt|jt|jtd|j DfS)Nc32K|]}t|ywrrqrhs rrkz2NVMMADistributedLayout.__hash__..6A#5:Ars)rtrrrrLrrPrs rruzNVMMADistributedLayout.__hash__4sNU4<<(%0B0B*CU4K[K[E\AAACD Drc,t|jSr)rTrLrs rrzNVMMADistributedLayout.rank8s4%%&&rrxr{s@rrrsr#Y9c"'"=JT#Y=Q  ^^D''rrcLeZdZdZedZdeedefdZ deedefdZ y) SharedLayoutz; Base class for shared memory layouts in Gluon IR. ct|Srrrs rrzSharedLayout.typeBrrrrc htj||Dcgc] }t|c}dScc}wrrrs rr!zSharedLayout.format_tensor_viewFr"r#c htj||Dcgc] }t|c}dScc}wr%rrs rr&z!SharedLayout.format_hardware_viewIr'r#N) r(r)r*r+r,rr-r.r/r!r&r0rrrr=sO$$_S _c_^$s)^^rrc|s|St|}t|d}dg|z}|D]7}t||k(sJt|D]}t||||||<9t|D]}||dk(rd||<||xxdzcc<t|D]-}||||zdk(s Jd|d|||xx||zcc</|S)NrrrzShape z is not divisible by CGA layout )r-rTrangemax)rrP shape_per_ctar cga_shaperir~s r_get_shape_per_ctarMs  KM z!} Dd I75zT!!!t 7Ay|U1X6IaL 77 4[ Q<1 IaL aLA L  T{.S!IcN2a7u6%Hhisht9uu7cy~-. rceZdZUdZeed<eed<dZeed<dZeed<dZ eed<e e Z e e eed <fd Zd Zeedd ZdefdZdZxZS)NVMMASharedLayouta 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. cga_layout (Optional[List[List[int]]]): Bases describing CTA tiling. swizzle_byte_widthelement_bitwidthrrF transposed fp4_paddedrNrPct|dt|jt|dt|jt|dt|j t|dt|j |jxsg}|rt|d|jk(sJt|dt|jt|dt||jdvsJ|jd vsJ|j r-|jd k(sJd |jd k(sJd yy)NrrrrrrrP) @)rrrrz+fp4_padded only supports 128 byte swizzlingrz3fp4_padded is only supported for element_bitwidth=8) rRrSrrrrrrPrTr)rrPrUs rrVzNVMMASharedLayout.__post_init__xs2 02FtG^G^2_` .0DTEZEZ0[\ L*>t*OP L*>t*OP__* z!}%2 22 F$8$CD L*>z*JK$$777&&*:::: ??**c1 `3` `1((A- d/d d- rc|j|j|j|j|j|j |j Sr)get_nvmma_shared_layoutrrrrrPrr6s rr8zNVMMASharedLayout._to_irsC..  # #  ! ! OO OO OO II   rcx|rdnd}||n 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. rrNrrrrr)rrrrrrP)rrTprimitive_bitwidthr) block_shapedtyperrrPpacking_factorrrcontig_dim_sizecontig_dim_bytesrflatten_outer_dimsizes rget_default_forz!NVMMASharedLayout.get_default_fors+)a'1'9 ?QR]_i?j ; )!"- bq0AAM'+n<*U-E-EEJ s "'7#'='B!$   #(82(=(B!#   #(82(=(B!# !" !#2& &D  %  & { a #4q#8!"  1"55!!!   rrc |jr"djd|jDnd}d|jd|jd|jd|j d|d S)Nr]c3ZK|]#}djtt|%ywrfr^rhs rrkz+NVMMASharedLayout.mangle..rlrmr\NVMMA__NVMMA)rPr_rrrrrs rr;zNVMMASharedLayout.manglesrUYUdUdSXXQQQjl //0$2G2G1H$//IZZ[\`\k\k[llmnxmyyA Arc t|j|j|j|j|j |j r!td|j DfSdfS)Nc32K|]}t|ywrrqrhs rrkz-NVMMASharedLayout.__hash__..rrs)rtrrrrrrPrrrs rruzNVMMASharedLayout.__hash__slT,,d.C.CTYYPTP_P_aeapapEI__UAAA`a aZ^`a ar)FFN)r(r)r*r+r.ryrrboolrrr-rPrrVr8 staticmethodr rr/r;rurzr{s@rrrds D#MJJ"'"=JT#Y=e* # # JAAarrceZdZUdZeed<eed<eed<eeed<eeZ eeeed<fdZ d Z d e fd Z d ZxZS) SwizzledSharedLayoutaq 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. cga_layout (Optional[List[List[int]]]): Bases describing CTA tiling. rj per_phase max_phaserMrNrPcdt|dt|jt|dt|jt|dt|j t|dt|j t|dt|jy)NrjrrrMrP)rRrSrrjrrrMrPrs rrVz"SwizzledSharedLayout.__post_init__s E#7#AB K)=dnn)MN K)=dnn)MN G%9$**%EF L*>t*OPrc|j|j|j|j|j|j Sr)get_swizzled_shared_layoutrjrrrMrPr6s rr8zSwizzledSharedLayout._to_irs811 HH NN NN JJ OO   rrc d}|jr"djd|jDnd}d|jd|jd|jd||j d|d S)NcF|ydjtt|Sr[r^ras rrcz.SwizzledSharedLayout.mangle..stringifyrdrr]c3ZK|]#}djtt|%ywrfr^rhs rrkz.SwizzledSharedLayout.mangle..rlrmr\SSS__SSS)rPr_rjrrrM)rrcrPs rr;zSwizzledSharedLayout.manglesq ) VZUdUdSXXQQQjl dhhZq 0$..1A9TZZCXBYYZ[eZffjkkrc t|j|j|jt |j t d|j DfS)Nc32K|]}t|ywrrqrhs rrkz0SwizzledSharedLayout.__hash__..sOv_bPUVYPZOvrs)rtrjrrrrrMrPrs rruzSwizzledSharedLayout.__hash__sF XXt~~t~~uTZZ7H%OvfjfufuOvJv wy yr)r(r)r*r+r.ryrrr-rPrVr8r/r;rurzr{s@rrrsY  HNN 9"'"=JT#Y=Q llyrrceZdZUdZeeeed<eeeed<eeeed<eeed<fdZdZde fd Z d Z e e gfd Zd ZxZS) PaddedSharedLayouta Represents a layout for the access to shared memory. Compared to SwizzledSharedLayout, it combined padding and element reordering via linear transformation (e.g. row permutation) 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]] with an identity remapping, the shared memory will be [e0, e1, p0, e2, e3, p1, p2, p3, e4, e5, p4, e6, e7, p5, p6, p7, ...] Furthermore this encoding allows for a linear remapping from the 1-D shared memory offset to logical n-D tensor elements. The remapping is given in the form of linear bases mapping from offset to [dim0, dim1...dimN-1]. See LinearLayout.h for more details how linear layouts are applied to remap elements. Some concrete examples using `xN` and `yN` to mean the logical n-D tensor elements and `pN` to mean padding: After padding for shape = [8] with interval-padding list [[2, 2]], offset_bases = [[2], [1]] and cga_layout = []: [x0, x2, p0 p1, x1, x3] After padding for shape = [8, 4] with interval_padding_pairs = [[8, 1]], offset_bases = [[0, 1], [0, 2], /*gap, stride by 2 rows*/[2, 0], [4, 0], [1, 0]]] and cga_layout = []: [ x0y0, x0y1, x0y2, x0y3, x2y0, x2y1, x2y2, x2y3, p0, x4y0, x4y1, x4y2, x4y3, x6y0, x6y1, x6y2, x6y3, p1, x1y0, x1y1, x1y2, x1y3, x3y0, x3y1, x3y2, x3y3, p2, x5y0, x5y1, x5y2, x5y3, x7y0, x7y1, x7y2, x7y3, ] Args: interval_padding_pairs (List[int]): List of [interval, padding] pair and both interval and padding must be powers of 2. offset_bases (List[int]): Bases for shared memory offsets cga_layout (List[List[int]]): Bases for block-level shared memory offsets. shape (List[int]): n-D logical shared memory shape interval_padding_pairs offset_basesrPrct|dt|jt|dt|jt|dt|j t|dt|j t|j }|jD]}t||k(rJ|j D]}t||k(rJ|jy)NrrrPr) rRrSrrrrPrrTverifyrs rrVz PaddedSharedLayout.__post_init__8s 4mDD_D_6`a NM$:K:K,LM L-*HI G]4::%>?4::&& &Eu:% %% &__ &Eu:% %% & rct|j\}}|j|||j|j|j Sr)ziprget_padded_shared_layoutrrPr)rr7 intervalspaddingss rr8zPaddedSharedLayout._to_irGsD!4#>#>? 8// 8TEVEVX\XgXgimisisttrrc pd|jd|jd|jd|jd S)N PaddedShared_r] _PaddedShared)rrrPrrs rr;zPaddedSharedLayout.mangleKsCt::;1T=N=N.Qs4d3t9>4sc&|dkDxr ||dz zdk(Srr0ns rz+PaddedSharedLayout.verify..W!a%".Xs7=#7z;PaddedSharedLayout interval values must all be power of twoc3.K|] }|ywrr0rs rrkz,PaddedSharedLayout.verify..Ys6=#6rz:PaddedSharedLayout padding values must all be power of twoz*PaddedSharedLayout order must not be empty)rrTrrr-setr)rpairsrrunique_intervalsrrs @rrzPaddedSharedLayout.verifyNs++5zA~www~4e4444!5k 8I/#$I666< 7Y77v9vv76X66t8tt64::axEEExrc  t|t|k(sJd t fd|DsJt|}|r t||n|}g}|D]c}tt t j ||D]4}|jt|D cgc]} | |k(rd|zndc} 6et||||Scc} w)zReturns a PaddedSharedLayout with the given interval and padding pairs and an identity mapping as the linear component for the given shape and order. c&|dkDxr ||dz zdk(Srr0rs rrz6PaddedSharedLayout.with_identity_for..drrc3.K|] }|ywrr0rs rrkz7PaddedSharedLayout.with_identity_for..es3=#3rrr) rTrrrr.mathlog2rr) rrrMrPrrrr~rrrs @rwith_identity_forz$PaddedSharedLayout.with_identity_for^s 5zSZ'''< 3U33335zAK*5*=QV   ZCs499]3-?#@AB Z##ERVK$Xq18Q%Z%B$XY Z Z""8, TYZZ%YsC c  tttt|jttt|jttt|j t|j fSr)rtrrr`rrrPrrs rruzPaddedSharedLayout.__hash__qs^U3ud&A&ABCU3uVZVgVgKhEi3udoo67tzz9JLM Mr)r(r)r*r+rr.ryrVr8r/r;rrr r!rurzr{s@rrrs8r!cO+tCy/!T#Y 9 u}}F KM[["MrrceZdZUdZeeeed<eeZ eeeed<dZ eed<fdZ dZ d e fd Zed Zd ZxZS) SharedLinearLayoutzGRepresents a shared memory layout defined via an explicit LinearLayout.rrNrr alignmentcLt|dt|jt|dt|jt|dt |j t|jdk7sJdt|jd}|dkDsJd|jD]}t||k(rJ|jD]}t||k(rJ|j dkDr|j |j dz zdk(sJdy)Nrrr%rz1SharedLinearLayout offset_bases must not be emptyrzKd3i>Is Kedd++ rr$cdg|z}|s|Sd}|D]D}tdt|Dd}||}||xxdzcc<1|r4|J||xxdzcc<F|S)Nrc32K|]\}}|dk7s |yw)rNr0)rirvs rrkz bases_per_dim..s=$!Qa1fA=s r)nextr/)basesrskip_broadcastresult non_zero_idxridxs r bases_per_dimr=s~S4ZF  L &=)E"2=tD ?L 3K1 K+ ++ < A %  & Mrct|trt|jt |St|t t frt|j|S|jSr) isinstancerr=rrTr}rrLr)layoutrs rrLrLsQ&12V..E ;; F[*:; <V]]E22###r)T) dataclassesrrr-rtypingrtriton.language.corerrrtriton.runtime.jitr triton._C.libtritonr r r3rBrIr}rrrrrrrrr$r=rLr0rrrFs( TT1(^^( $ 3" 3 3 $ 8' 8 8 $:%::z $'X#'X'XT $6/66r $0(00f $4 ''.''!''T ^ ^ , $_a _a_aD $4 -y<-y!-y` $4 zMzM!zMz $* * * \,$r