o iX@srddlmZddlmZmZddlmZmZmZddl m Z ddZ GdddZ ed d Gd d d e Z ed d Gd dde Zed d Gddde Zed d Gddde Zed d Gddde Zed d dGddde ZGdddZe ddZed d GdddeZed d dGdddeZed d dGd d!d!eZd'd"d#Zd$d%Zd&S)() dataclass)ListOptional)_unwrap_if_constexpr _unwrap_shapeconstexpr_type)constexpr_functioncCsd|jpdg|}|jpdg|}|jpttt|}t|d|t|d|t|d|dS)N ctas_per_cga cta_split_num cta_order)r r r listreversedrangeobject __setattr__)layoutrankr r r ra/home/ubuntu/vllm_env/lib/python3.10/site-packages/triton/experimental/gluon/language/_layouts.py_realize_cta_layouts rc@eZdZdZeddZdS)DistributedLayoutz@ Base class for distributed memory layouts in Gluon IR. cCt|SNrselfrrrtypezDistributedLayout.typeN__name__ __module__ __qualname____doc__propertyrrrrrrrT)frozenc@seZdZddZddZdS) AutoLayoutcCs|Sr)get_auto_layoutrbuilderrrr_to_irszAutoLayout._to_ircCsdS)NALrrrrrmangle szAutoLayout.mangleN)r!r"r#r,r.rrrrr(s r(cseZdZUdZeeed<eeed<eeed<eeed<dZeeeed<dZ eeeed<dZ eeeed <fd d Z d d Z de fddZddZZS) 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 r cstdt|jtdt|jtdt|jtdt|jtdt|jtdt|jtdt|j t |j}t ||t |j|ksYJt |j|ksbJt |j|kskJt |j|kstJt |j|ks}Jt |j |ksJdS)Nr0r1r2r3r r r ) superrrr0r1r2r3r r r lenrrr __class__rr __post_init__:s  zBlockedLayout.__post_init__c C$||j|j|j|j|j|j|jSr)get_blocked_layoutr0r1r2r3r r r r*rrrr,LzBlockedLayout._to_irreturnc Cs~dd}||j}||j}||j}||j}||j}||j}||j}d|d|d|d|d|d|d|dS)NcS|durdSdtt|SN_joinmapstrxrrr stringifyYz'BlockedLayout.mangle..stringifyB)r0r1r2r3r r r ) rrHr0r1r2r3r r r rrrr.Ws       0zBlockedLayout.manglec Csbtt|jt|jt|jt|j|jrt|jnd|jr"t|jnd|jr-t|jfSdfSr) hashtupler0r1r2r3r r r rrrr__hash__gszBlockedLayout.__hash__r!r"r#r$rint__annotations__r rr r r9r,rEr.rM __classcell__rrr7rr/$s      r/csPeZdZUdZeed<eed<fddZddZde fd d Z d d Z Z 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. dimparentcs,tdt|jtdt|jdS)NrSrT)r4rrrSrTrr7rrr9szSliceLayout.__post_init__cCs||j|j|Sr)get_slice_layoutrSrTr,r*rrrr,s zSliceLayout._to_irr=cCsd|jd|jdS)NSLrA)rSrTr.rrrrr.zSliceLayout.manglecCst|j|jfSr)rKrSrTrrrrrMszSliceLayout.__hash__ r!r"r#r$rOrPrr9r,rEr.rMrQrrr7rrRss  rRcseZdZUdZeeeed<eeeed<eeeed<eeeed<eeed<fddZd d Zd d Z d dZ Z 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_basesshapecstdt|jtdt|jtdt|jtdt|jtdt|jt|j}|jD] }t||ksDJq:|jD] }t||ksRJqH|jD] }t||ks`JqV|jD] }t||ksnJqddS)NrZr[r\r]r^) r4rrrZr[r\r]r^r5)rrbasisr7rrr9s     z%DistributedLinearLayout.__post_init__cCs||j|j|j|j|jSr)get_distributed_linear_layoutrZr[r\r]r^r*rrrr,szDistributedLinearLayout._to_irc Cs.d|jd|jd|jd|jd|jd S)NDLLrA)rZr[r\r]r^rrrrr.s.zDistributedLinearLayout.manglecCsHtttt|jttt|jttt|jttt|jt|jfSr)rKrLrDrZr[r\r]r^rrrrrMsz DistributedLinearLayout.__hash__) r!r"r#r$rrOrPr9r,r.rMrQrrr7rrYs    rYcsXeZdZUdZeed<eed<eed<fddZddZd e fd d Z d d Z Z 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_indexrTk_widthcs@tdt|jtdt|jtdt|jdS)NrcrTrd)r4rrrcrTrdrr7rrr9szDotOperandLayout.__post_init__cCs||j|j||jSr)get_dot_operand_layoutrcrTr,rdr*rrrr,rWzDotOperandLayout._to_irr=cCs"d|jd|jd|jdS)NDOrA)rcrTr.rdrrrrr.s"zDotOperandLayout.manglecCst|j|j|jfSr)rKrcrTrdrrrrrMszDotOperandLayout.__hash__rXrrr7rrbs  rb)r'eqcseZdZUdZeeed<eeed<eeed<dZeeeed<dZ eeeed<dZ eeeed<fd d Z d d Z d e fddZddZZS)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. versionr2 instr_shapeNr r r cstdt|jtdt|jtdt|jtdt|jtdt|jtdt|jt |j}t ||t |j|ksOJt |j|ksXJt |j|ksaJdS)Nrir2rjr r r ) r4rrrir2rjr r r r5rr6r7rrr9s  z$NVMMADistributedLayout.__post_init__cCs ||j|j|j|j|j|jSr)get_mma_layoutrir2r r r rjr*rrrr,szNVMMADistributedLayout._to_irr=c Cs6d|jd|jd|jd|jd|jd|jd S)NMMA_rA_MMA)rir2rjr r r rrrrr. s6zNVMMADistributedLayout.manglecCsZtt|jt|jt|j|jrt|jnd|jrt|jnd|jr)t|jfSdfSr)rKrLrir2rjr r r rrrrrM szNVMMADistributedLayout.__hash__rNrrr7rrhs    rhc@r) SharedLayoutz; Base class for shared memory layouts in Gluon IR. cCrrrrrrrrrzSharedLayout.typeNr rrrrrnr&rncCsJ|}|dur#t|t|ksJtt|D] }||||<q|Sr)r5r)r^r shape_per_ctarSrrr_get_shape_per_ctas rpcseZdZUdZeed<eed<eed<dZeed<dZeed<dZ e e eed <dZ e e eed <dZ e e eed <fd d ZddZee  dddZdefddZddZZS)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 r cstdt|jtdt|jtdt|jtdt|jtdt|jtdt|jtdt|j tdt|j |jd vsWJ|jd vs^J|j}t ||t |j|ksoJt |j |ksxJt |j |ksJdS) Nrrrsrrtrur r r ) @)rrxry) r4rrrrrsrrtrur r r rr5r6r7rrr9?s zNVMMASharedLayout.__post_init__c Cr:r)get_nvmma_shared_layoutrrrsrtrur r r r*rrrr,Qr<zNVMMASharedLayout._to_irc Cs|rdnd}t||}t|} |r|dd|dd}|d|} | |jd} | dkr7| ddkr7d} n| dkrD| ddkrDd} n| d krQ| d dkrQd } nd} d} |ddD]}| |9} q[t|dksl| dkrnd} t| |j| |||||d 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 Nrvrzrryrx)rrrsrrtrur r r )rpr5primitive_bitwidthrq) block_shapedtypertrur r r packing_factorrorcontig_dim_sizecontig_dim_bytesrrflatten_outer_dimsizerrrget_default_for\s:   z!NVMMASharedLayout.get_default_forr=c Cs&d|jd|jd|jd|jd S)NNVMMA_rA_NVMMA)rrrsrtrurrrrr.s&zNVMMASharedLayout.manglec CsVt|j|j|j|j|j|jrt|jnd|jrt|jnd|j r't|j fSdfSr) rKrrrsrrtrur rLr r rrrrrMszNVMMASharedLayout.__hash__)FFNNN)r!r"r#r$rOrPrtboolrur rrr r r9r, staticmethodrrrEr.rMrQrrr7rrq's&      (rqcseZdZUdZeed<eed<eed<eeed<dZeeeed<dZ eeeed<dZ eeeed <fd d Z d d Z de fddZddZZS)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_phaser3Nr r r cstdt|jtdt|jtdt|jtdt|jtdt|jtdt|jtdt|j t |j}t ||t |j|ksYJt |j|ksbJt |j |kskJdS)Nrrrr3r r r ) r4rrrrrr3r r r r5rr6r7rrr9s  z"SwizzledSharedLayout.__post_init__c Cr:r)get_swizzled_shared_layoutrrrr3r r r r*rrrr,r<zSwizzledSharedLayout._to_irr=cCsVdd}d|jd|jd|jd||jd||jd||jd||jdS)NcSr>r?rBrFrrrrHrIz.SwizzledSharedLayout.mangle..stringifySSS_rA_SSS)rrrr3r r r rrHrrrr.sNzSwizzledSharedLayout.manglec CsVt|j|j|jt|j|jrt|jnd|jrt|jnd|jr't|jfSdfSr) rKrrrrLr3r r r rrrrrMszSwizzledSharedLayout.__hash__)r!r"r#r$rOrPrr rr r r9r,rEr.rMrQrrr7rrs      rcseZdZUdZeeeed<eeed<dZeeeed<dZ eeeed<dZ eeeed<fdd Z d d Z d e fd dZddZddZZS)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_pairsr3Nr r r csptdt|jtdt|jtdt|jtdt|jtdt|j| dS)Nrr3r r r ) r4rrrrr3r r r verifyrr7rrr9s  z PaddedSharedLayout.__post_init__cCs*t|j\}}||||j|j|j|jSr)ziprget_padded_shared_layoutr3r r r )rr+ intervalspaddingsrrrr,szPaddedSharedLayout._to_irr=c CsJdd}d||jd||jd||jd||jd||jd S)NcSr>r?rBrFrrrrHrIz,PaddedSharedLayout.mangle..stringify PaddedShared_rA _PaddedShared)rr3r r r rrrrr.sBzPaddedSharedLayout.manglecs|j}t|dks Jdtdd|DsJt|\}}tt|}t|t|ks.Jddtfdd|DsAJdtfd d|DsPJd t|j}|dks]Jd t||t|j|kskJt|j |kstJt|j |ks}JdS) NrzVPaddedSharedLayout interval_padding_pairs must have at least one interval-padding paircss|] }t|dkVqdS)r|N)r5).0pairrrr sz,PaddedSharedLayout.verify..cSs|dko ||d@dkS)Nrr r)nrrrsz+PaddedSharedLayout.verify..c3|]}|VqdSrrrr is_power_of_2rrrz;PaddedSharedLayout interval values must all be power of twoc3rrrrrrrrrz:PaddedSharedLayout padding values must all be power of twoz*PaddedSharedLayout order must not be empty) rr5allrr setr3rr r r )rpairsrrunique_intervalsrrrrrs    zPaddedSharedLayout.verifycCsXtttt|jt|j|jrt|jnd|jrt|jnd|jr(t|jfSdfSr)rKrLrDrr3r r r rrrrrM#szPaddedSharedLayout.__hash__)r!r"r#r$rrOrPr rr r r9r,rEr.rrMrQrrr7rrs     rcCszdg|}|s |Sd}|D]-}tddt|Dd}|dur*|}||d9<q |s:|dus2J||d9<q |S)Nr css |] \}}|dkr|VqdS)rNr)rivrrrr5sz bases_per_dim..r|)next enumerate)basesrskip_broadcastresult non_zero_idxr_idxrrr bases_per_dim+s  rcCs:t|tr t|jt|St|ttfrt|j|S|jSr) isinstancerYrr\r5rRrbr2rT)rr^rrrr2As  r2N)T) dataclassesrtypingrrtriton.language.corerrrtriton.runtime.jitrrrr(r/rRrYrbrhrnrprqrrrr2rrrrs:     N5 /   h ? Z