§
    OŠtjZe  ã                   ó|  — d dl mZmZ d dlZd dlZd dlmZ d dlmZm	Z	m
Z
 d dlmZ d dlmZ  G d„ d¦  «        Z ed	¬
¦  «         G d„ de¦  «        ¦   «         Z ed	¬
¦  «         G d„ de¦  «        ¦   «         Z ed	¬
¦  «         G d„ de¦  «        ¦   «         Z ed	¬
¦  «         G d„ de¦  «        ¦   «         Z ed	¬
¦  «         G d„ de¦  «        ¦   «         Z ed	¬
¦  «         G d„ de¦  «        ¦   «         Z ed	d	¬¦  «         G d„ de¦  «        ¦   «         Z G d„ d¦  «        Zed„ ¦   «         Z ed	¬
¦  «         G d„ de¦  «        ¦   «         Z ed	d	¬¦  «         G d„ d e¦  «        ¦   «         Z ed	d	¬¦  «         G d!„ d"e¦  «        ¦   «         Z ed	¬
¦  «         G d#„ d$e¦  «        ¦   «         Zd'd%„Zd&„ ZdS )(é    )Ú	dataclassÚfieldN)ÚList)Ú_unwrap_if_constexprÚ_unwrap_shapeÚconstexpr_type)Úconstexpr_function)Úgluon_irc                   óv   — e Zd ZdZed„ ¦   «         Zed„ ¦   «         Zdee         de	fd„Z
dee         de	fd„ZdS )	ÚDistributedLayoutz@
    Base class for distributed memory layouts in Gluon IR.
    c                 ó    — t          | ¦  «        S ©N©r   ©Úselfs    úi/var/www/html/CA-Chatbot/venv/lib/python3.11/site-packages/triton/experimental/gluon/language/_layouts.pyÚtypezDistributedLayout.type   ó   € å˜dÑ#Ô#Ð#ó    c                 ó    — t          d¦  «        ‚)Nz-DistributedLayout subclasses must define rank)ÚNotImplementedErrorr   s    r   ÚrankzDistributedLayout.rank   s   € å!Ð"QÑRÔRÐRr   ÚshapeÚreturnc                 óB   — t          j        | d„ |D ¦   «         d¦  «        S )Nc                 ó,   — g | ]}t          |¦  «        ‘ŒS © ©r   ©Ú.0Úss     r   ú
<listcomp>z8DistributedLayout.format_tensor_view.<locals>.<listcomp>   ó!   € Ð.VÐ.VÐ.VÈ1Õ/CÀAÑ/FÔ/FÐ.VÐ.VÐ.Vr   F©r
   Úget_layout_view©r   r   s     r   Úformat_tensor_viewz$DistributedLayout.format_tensor_view   ó(   € ÝÔ'¨Ð.VÐ.VÐPUÐ.VÑ.VÔ.VÐX]Ñ^Ô^Ð^r   c                 óB   — t          j        | d„ |D ¦   «         d¦  «        S )Nc                 ó,   — g | ]}t          |¦  «        ‘ŒS r   r   r   s     r   r"   z:DistributedLayout.format_hardware_view.<locals>.<listcomp>   r#   r   Tr$   r&   s     r   Úformat_hardware_viewz&DistributedLayout.format_hardware_view   ó(   € ÝÔ'¨Ð.VÐ.VÐPUÐ.VÑ.VÔ.VÐX\Ñ]Ô]Ð]r   N)Ú__name__Ú
__module__Ú__qualname__Ú__doc__Úpropertyr   r   ÚlistÚintÚstrr'   r+   r   r   r   r   r      s±   € € € € € ðð ð ð$ð $ñ „Xð$ð ðSð Sñ „XðSð_¨¨S¬	ð _°cð _ð _ð _ð _ð^¨$¨s¬)ð ^¸ð ^ð ^ð ^ð ^ð ^ð ^r   r   T)Úfrozenc                   ó0   — e Zd Zd„ Zd„ Zed„ ¦   «         ZdS )Ú
AutoLayoutc                 ó*   — |                      ¦   «         S r   )Úget_auto_layout©r   Úbuilders     r   Ú_to_irzAutoLayout._to_ir"   s   € Ø×&Ò&Ñ(Ô(Ð(r   c                 ó   — dS )NÚALr   r   s    r   ÚmanglezAutoLayout.mangle%   ó   € Øˆtr   c                 ó    — t          d¦  «        ‚)NzAutoLayout has no rank©Ú
ValueErrorr   s    r   r   zAutoLayout.rank(   s   € åÐ1Ñ2Ô2Ð2r   N©r-   r.   r/   r<   r?   r1   r   r   r   r   r7   r7      sM   € € € € € ð)ð )ð )ðð ð ð ð3ð 3ñ „Xð3ð 3ð 3r   r7   c                   ó0   — e Zd Zd„ Zd„ Zed„ ¦   «         ZdS )ÚCoalescedLayoutc                 ó*   — |                      ¦   «         S r   )Úget_coalesced_layoutr:   s     r   r<   zCoalescedLayout._to_ir0   s   € Ø×+Ò+Ñ-Ô-Ð-r   c                 ó   — dS )NÚCLr   r   s    r   r?   zCoalescedLayout.mangle3   r@   r   c                 ó    — t          d¦  «        ‚)NzCoalescedLayout has no rankrB   r   s    r   r   zCoalescedLayout.rank6   s   € åÐ6Ñ7Ô7Ð7r   NrD   r   r   r   rF   rF   -   sM   € € € € € ð.ð .ð .ðð ð ð ð8ð 8ñ „Xð8ð 8ð 8r   rF   c                   óä   ‡ — e Zd ZU dZee         ed<   ee         ed<   ee         ed<   ee         ed<    ee¬¦  «        Z	eee                  ed<   ˆ fd„Z
d	„ Zd
efd„Zd„ Zed„ ¦   «         Zˆ xZS )Ú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_threadÚthreads_per_warpÚwarps_per_ctaÚorder©Údefault_factoryÚ
cga_layoutc                 óÔ  •— t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          | j        ¦  «        }t          | j        ¦  «        |k    sJ ‚t          | j        ¦  «        |k    sJ ‚t          | j        ¦  «        |k    sJ ‚d S )NrN   rO   rP   rQ   rT   )	ÚsuperÚ__setattr__r   rN   rO   rP   rQ   rT   Úlen)r   r   Ú	__class__s     €r   Ú__post_init__zBlockedLayout.__post_init__M   s(  ø€ Ý‰Œ×ÒÐ-Õ/CÀDÔDXÑ/YÔ/YÑZÔZÐZÝ‰Œ×ÒÐ.Õ0DÀTÔEZÑ0[Ô0[Ñ\Ô\Ð\Ý‰Œ×Ò˜OÕ-AÀ$ÔBTÑ-UÔ-UÑVÔVÐVÝ‰Œ×Ò˜GÕ%9¸$¼*Ñ%EÔ%EÑFÔFÐFÝ‰Œ×Ò˜LÕ*>¸t¼Ñ*OÔ*OÑPÔPÐPå�4Ô'Ñ(Ô(ˆÝ�4Ô(Ñ)Ô)¨TÒ1Ð1Ð1Ð1Ý�4Ô%Ñ&Ô&¨$Ò.Ð.Ð.Ð.Ý�4”:‰Œ $Ò&Ð&Ð&Ð&Ð&Ð&r   c                 óf   — |                      | j        | j        | j        | j        | j        ¦  «        S r   )Úget_blocked_layoutrN   rO   rP   rQ   rT   r:   s     r   r<   zBlockedLayout._to_irY   s6   € Ø×)Ò)ØÔ ØÔ!ØÔØŒJØŒOñ
ô 
ð 	
r   r   c                 ó  — d„ } || j         ¦  «        } || j        ¦  «        } || j        ¦  «        } || j        ¦  «        }| j        r$d                     d„ | j        D ¦   «         ¦  «        nd}d|› d|› d|› d|› d|› d�S )Nc                 óZ   — | €dS d                      t          t          | ¦  «        ¦  «        S ©NÚ Ú_©ÚjoinÚmapr4   ©Úxs    r   Ú	stringifyz'BlockedLayout.mangle.<locals>.stringifyd   ó&   € ØˆyØ�rØ—8’8�C¥ Q™KœKÑ(Ô(Ð(r   ra   c              3   óf   K  — | ],}d                       t          t          |¦  «        ¦  «        V — Œ-dS ©ú~Nrb   ©r    Úvecs     r   ú	<genexpr>z'BlockedLayout.mangle.<locals>.<genexpr>m   ó6   è è € ÐQÐQ¸#˜cŸhšh¥s­3°¡}¤}Ñ5Ô5ÐQÐQÐQÐQÐQÐQr   r`   ÚB)rN   rO   rP   rQ   rT   rc   )r   rg   rN   rO   rP   rQ   rT   s          r   r?   zBlockedLayout.mangleb   s¾   € ð	)ð 	)ð 	)ð
 $˜) DÔ$8Ñ9Ô9ˆØ$˜9 TÔ%:Ñ;Ô;ÐØ!˜	 $Ô"4Ñ5Ô5ˆØ�	˜$œ*Ñ%Ô%ˆØUYÔUdÐl�S—X’XÐQÐQÀÄÐQÑQÔQÑQÔQÐQÐjlˆ
Ø]�?Ð]Ð]Ð%5Ð]Ð]¸Ð]Ð]ÈÐ]Ð]ÐPZÐ]Ð]Ð]Ð]r   c           
      óò   — t          t          | j        ¦  «        t          | j        ¦  «        t          | j        ¦  «        t          | j        ¦  «        t          d„ | j        D ¦   «         ¦  «        f¦  «        S )Nc              3   ó4   K  — | ]}t          |¦  «        V — Œd S r   ©Útuplerl   s     r   rn   z)BlockedLayout.__hash__.<locals>.<genexpr>r   s(   è è € Ð-TÐ-T¸S­e°C©j¬jÐ-TÐ-TÐ-TÐ-TÐ-TÐ-Tr   )Úhashrt   rN   rO   rP   rQ   rT   r   s    r   Ú__hash__zBlockedLayout.__hash__p   sn   € Ý•U˜4Ô/Ñ0Ô0µ%¸Ô8MÑ2NÔ2NÕPUÐVZÔVhÑPiÔPiÝ˜4œ:Ñ&Ô&­Ð-TÐ-TÀDÄOÐ-TÑ-TÔ-TÑ(TÔ(TðVñ Wô Wð 	Wr   c                 ó*   — t          | j        ¦  «        S r   )rX   rQ   r   s    r   r   zBlockedLayout.rankt   ó   € å�4”:‰ŒÐr   ©r-   r.   r/   r0   r   r3   Ú__annotations__r   r2   rT   rZ   r<   r4   r?   rv   r1   r   Ú__classcell__©rY   s   @r   rM   rM   ;   s  ø€ € € € € € ð	ð 	ð ˜#”YÐÐÑØ˜3”iÐÐÑØ˜”9ÐÐÑØ�Œ9ÐÐÑØ"' %¸Ð"=Ñ"=Ô"=€J��T˜#”Y”Ð=Ð=Ñ=ð
'ð 
'ð 
'ð 
'ð 
'ð
ð 
ð 
ð^˜ð ^ð ^ð ^ð ^ðWð Wð Wð ðð ñ „Xðð ð ð ð r   rM   c                   ó|   ‡ — e Zd ZU dZeed<   eed<   ˆ fd„Zd„ Zde	fd„Z
d„ Zed	„ ¦   «         Zed
„ ¦   «         Zˆ x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.
    ÚdimÚparentc                 óØ   •— t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         d S )Nr   r€   )rV   rW   r   r   r€   ©r   rY   s    €r   rZ   zSliceLayout.__post_init__…   sS   ø€ Ý‰Œ×Ò˜EÕ#7¸¼Ñ#AÔ#AÑBÔBÐBÝ‰Œ×Ò˜HÕ&:¸4¼;Ñ&GÔ&GÑHÔHÐHÐHÐHr   c                 óh   — |                      | j        | j                             |¦  «        ¦  «        S r   )Úget_slice_layoutr   r€   r<   r:   s     r   r<   zSliceLayout._to_ir‰   s2   € Ø×'Ò'ØŒHØŒK×Ò˜wÑ'Ô'ñ
ô 
ð 	
r   r   c                 óL   — d| j         › d| j                             ¦   «         › d�S )NÚSLra   )r   r€   r?   r   s    r   r?   zSliceLayout.mangle�   s+   € Ø7�D”HÐ7Ð7˜tœ{×1Ò1Ñ3Ô3Ð7Ð7Ð7Ð7r   c                 ó8   — t          | j        | j        f¦  «        S r   )ru   r   r€   r   s    r   rv   zSliceLayout.__hash__’   s   € Ý�T”X˜tœ{Ð+Ñ,Ô,Ð,r   c                 ó    — | j         j        dz
  S ©Né   ©r€   r   r   s    r   r   zSliceLayout.rank•   s   € àŒ{Ô !Ñ#Ð#r   c                 ó„   ‡ — ‰ j         j        }|sg S ‰ j         j        }d‰ j        cxk    r|k     sn J ‚ˆ fd„|D ¦   «         S )Nr   c                 óT   •— g | ]$}|d ‰j         …         |‰j         dz   d …         z   ‘Œ%S r‰   )r   )r    Úbasisr   s     €r   r"   z*SliceLayout.cga_layout.<locals>.<listcomp>¡   s8   ø€ ÐWÐWÐW¸E��i�t”x�iÔ  5¨¬°A©¨¨Ô#7Ñ7ÐWÐWÐWr   )r€   rT   r   r   )r   Úparent_cga_layoutr   s   `  r   rT   zSliceLayout.cga_layout™   sg   ø€ à œKÔ2ÐØ ð 	ØˆIàŒ{ÔˆØ�D”HÐ#Ð#Ò#Ð#˜tÒ#Ð#Ð#Ð#Ð#Ð#ØWÐWÐWÐWÐEVÐWÑWÔWÐWr   ©r-   r.   r/   r0   r3   rz   r   rZ   r<   r4   r?   rv   r1   r   rT   r{   r|   s   @r   r~   r~   y   sÖ   ø€ € € € € € ðð ð 
€H€H�HØÐÐÑðIð Ið Ið Ið Ið
ð 
ð 
ð8˜ð 8ð 8ð 8ð 8ð-ð -ð -ð ð$ð $ñ „Xð$ð ðXð Xñ „XðXð Xð Xð Xð Xr   r~   c                   óê   ‡ — e Zd ZU dZeee                  ed<   eee                  ed<   eee                  ed<   eee                  ed<   ee         ed<   ˆ fd„Zd„ Zd	„ Z	d
„ Z
ed„ ¦   «         Zˆ x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_basesr   c                 ó0  •— t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          | j        ¦  «        }| j        D ]}t          |¦  «        |k    sJ ‚Œ| j        D ]}t          |¦  «        |k    sJ ‚Œ| j        D ]}t          |¦  «        |k    sJ ‚Œ| j        D ]}t          |¦  «        |k    sJ ‚Œd S )Nr“   r”   r•   r–   r   )	rV   rW   r   r“   r”   r•   r–   r   rX   ©r   r   rŽ   rY   s      €r   rZ   z%DistributedLinearLayout.__post_init__·   s|  ø€ Ý‰Œ×Ò˜K­°t´~Ñ)FÔ)FÑGÔGÐGÝ‰Œ×Ò˜L­-¸¼Ñ*HÔ*HÑIÔIÐIÝ‰Œ×Ò˜L­-¸¼Ñ*HÔ*HÑIÔIÐIÝ‰Œ×Ò˜M­=¸Ô9IÑ+JÔ+JÑKÔKÐKÝ‰Œ×Ò˜G¥]°4´:Ñ%>Ô%>Ñ?Ô?Ð?å�4”:‰Œˆà”^ð 	&ð 	&ˆEÝ�u‘:”: Ò%Ð%Ð%Ð%Ð%Ø”_ð 	&ð 	&ˆEÝ�u‘:”: Ò%Ð%Ð%Ð%Ð%Ø”_ð 	&ð 	&ˆEÝ�u‘:”: Ò%Ð%Ð%Ð%Ð%ØÔ%ð 	&ð 	&ˆEÝ�u‘:”: Ò%Ð%Ð%Ð%Ð%ð	&ð 	&r   c                 óf   — |                      | j        | j        | j        | j        | j        ¦  «        S r   )Úget_distributed_linear_layoutr“   r”   r•   r–   r   r:   s     r   r<   zDistributedLinearLayout._to_irÉ   s6   € Ø×4Ò4°T´^ÀTÄ_ÐVZÔVeÐgkÔgwØ59´ZñAô Að 	Ar   c                 óX   — d| j         › d| j        › d| j        › d| j        › d| j        › d�S )NÚDLLra   )r“   r”   r•   r–   r   r   s    r   r?   zDistributedLinearLayout.mangleÍ   s@   € Øl�T”^ÐlÐl d¤oÐlÐl¸¼ÐlÐlÈ$ÔJZÐlÐlÐ]aÔ]gÐlÐlÐlÐlr   c                 óv  — t          t          t          t          | j        ¦  «        ¦  «        t          t          t          | j        ¦  «        ¦  «        t          t          t          | j        ¦  «        ¦  «        t          t          t          | j        ¦  «        ¦  «        t          | j        ¦  «        f¦  «        S r   )ru   rt   rd   r“   r”   r•   r–   r   r   s    r   rv   z DistributedLinearLayout.__hash__Ð   sƒ   € ÝÝ•#•e˜Tœ^Ñ,Ô,Ñ-Ô-Ý•#•e˜Tœ_Ñ-Ô-Ñ.Ô.Ý•#•e˜Tœ_Ñ-Ô-Ñ.Ô.Ý•#•e˜TÔ-Ñ.Ô.Ñ/Ô/Ý�$”*ÑÔð
ñ ô ð 	r   c                 ó*   — t          | j        ¦  «        S r   )rX   r   r   s    r   r   zDistributedLinearLayout.rankÙ   rx   r   )r-   r.   r/   r0   r   r3   rz   rZ   r<   r?   rv   r1   r   r{   r|   s   @r   r’   r’   ¤   sô   ø€ € € € € € ð
ð 
ð �D˜”IŒÐÐÑØ�T˜#”Y”ÐÐÑØ�T˜#”Y”ÐÐÑØ�d˜3”i”Ð Ð Ñ Ø�Œ9ÐÐÑð&ð &ð &ð &ð &ð$Að Að Aðmð mð mðð ð ð ðð ñ „Xðð ð ð ð r   r’   c                   ó†   ‡ — e Zd ZU dZeed<   eed<   eed<   ˆ fd„Zd„ Zde	fd„Z
d	„ Zed
„ ¦   «         Zed„ ¦   «         Zˆ x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_indexr€   Úk_widthc                 ó@  •— t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         d S )Nr¡   r€   r¢   )rV   rW   r   r¡   r€   r¢   r‚   s    €r   rZ   zDotOperandLayout.__post_init__ì   sy   ø€ Ý‰Œ×Ò˜OÕ-AÀ$ÔBTÑ-UÔ-UÑVÔVÐVÝ‰Œ×Ò˜HÕ&:¸4¼;Ñ&GÔ&GÑHÔHÐHÝ‰Œ×Ò˜IÕ';¸D¼LÑ'IÔ'IÑJÔJÐJÐJÐJr   c                 ót   — |                      | j        | j                             |¦  «        | j        ¦  «        S r   )Úget_dot_operand_layoutr¡   r€   r<   r¢   r:   s     r   r<   zDotOperandLayout._to_irñ   s2   € Ø×-Ò-¨dÔ.@À$Ä+×BTÒBTÐU\ÑB]ÔB]Ð_cÔ_kÑlÔlÐlr   r   c                 ó\   — d| j         › d| j                             ¦   «         › d| j        › d�S )NÚDOra   )r¡   r€   r?   r¢   r   s    r   r?   zDotOperandLayout.mangleô   s6   € ØP�DÔ&ÐPÐP¨¬×);Ò);Ñ)=Ô)=ÐPÐPÀÄÐPÐPÐPÐPr   c                 óD   — t          | j        | j        | j        f¦  «        S r   )ru   r¡   r€   r¢   r   s    r   rv   zDotOperandLayout.__hash__÷   s   € Ý�TÔ'¨¬°d´lÐCÑDÔDÐDr   c                 ó   — | j         j        S r   r‹   r   s    r   r   zDotOperandLayout.rankú   s   € àŒ{ÔÐr   c                 óX  ‡— t          t          | j        dg ¦  «        ¦  «        pg }|sg S | j        j        Št	          ˆfd„|D ¦   «         ¦  «        sJ ‚| j        dk    r‰dz
  n‰dz
  }d|cxk    r‰k     sn J ‚g }|D ]+}t          |¦  «        }d||<   |                     |¦  «         Œ,|S )NrT   c              3   ó>   •K  — | ]}t          |¦  «        ‰k    V — Œd S r   ©rX   )r    rŽ   r   s     €r   rn   z.DotOperandLayout.cga_layout.<locals>.<genexpr>  s.   øè è € ÐEÐE¨%•3�u‘:”: Ò%ÐEÐEÐEÐEÐEÐEr   r   rŠ   é   )r   Úgetattrr€   r   Úallr¡   r2   Úappend)r   r�   Úk_dimÚderivedrŽ   Ú	new_basisr   s         @r   rT   zDotOperandLayout.cga_layoutþ   sí   ø€ å0µ¸¼ÀlÐTVÑ1WÔ1WÑXÔXÐ^Ð\^ÐØ ð 	ØˆIàŒ{ÔˆÝÐEÐEÐEÐEÐ3DÐEÑEÔEÑEÔEÐEÐEÐEà Ô.°!Ò3Ð3��q‘�¸À¹ˆØ�EÐ Ð Ò Ð ˜DÒ Ð Ð Ð Ð Ð àˆØ&ð 	&ð 	&ˆEÝ˜U™œˆIØ ˆI�eÑØ�NŠN˜9Ñ%Ô%Ð%Ð%Øˆr   r�   r|   s   @r   r    r    Þ   sæ   ø€ € € € € € ðð ð ÐÐÑØÐÐÑØ€L€L�LðKð Kð Kð Kð Kð
mð mð mðQ˜ð Qð Qð Qð QðEð Eð Eð ð ð  ñ „Xð ð ðð ñ „Xðð ð ð ð r   r    )r5   Úeqc                   óÎ   ‡ — e Zd ZU dZee         ed<   ee         ed<   ee         ed<    ee¬¦  «        Z	eee                  ed<   ˆ fd„Z
d„ Zd	efd
„Zd„ Zed„ ¦   «         Zˆ xZS )Ú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.
    ÚversionrP   Úinstr_shaperR   rT   c                 ó¨  •— t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         d S )Nr·   rP   r¸   rT   )rV   rW   r   r·   rP   r¸   rT   r‚   s    €r   rZ   z$NVMMADistributedLayout.__post_init__"  sŸ   ø€ Ý‰Œ×Ò˜IÕ';¸D¼LÑ'IÔ'IÑJÔJÐJÝ‰Œ×Ò˜OÕ-AÀ$ÔBTÑ-UÔ-UÑVÔVÐVÝ‰Œ×Ò˜MÕ+?ÀÔ@PÑ+QÔ+QÑRÔRÐRÝ‰Œ×Ò˜LÕ*>¸t¼Ñ*OÔ*OÑPÔPÐPÐPÐPr   c                 óZ   — |                      | j        | j        | j        | j        ¦  «        S r   )Úget_mma_layoutr·   rP   rT   r¸   r:   s     r   r<   zNVMMADistributedLayout._to_ir(  s0   € Ø×%Ò%ØŒLØÔØŒOØÔñ	
ô 
ð 	
r   r   c           	      ó˜   — | j         r$d                     d„ | j         D ¦   «         ¦  «        nd}d| j        › d| j        › d| j        › d|› d�	S )Nra   c              3   óf   K  — | ],}d                       t          t          |¦  «        ¦  «        V — Œ-dS rj   rb   rl   s     r   rn   z0NVMMADistributedLayout.mangle.<locals>.<genexpr>1  ro   r   r`   ÚMMA_Ú_MMA)rT   rc   r·   rP   r¸   ©r   rT   s     r   r?   zNVMMADistributedLayout.mangle0  sc   € ØUYÔUdÐl�S—X’XÐQÐQÀÄÐQÑQÔQÑQÔQÐQÐjlˆ
Ø]�d”lÐ]Ð] TÔ%7Ð]Ð]¸$Ô:JÐ]Ð]ÈZÐ]Ð]Ð]Ð]r   c           	      óÌ   — t          t          | j        ¦  «        t          | j        ¦  «        t          | j        ¦  «        t          d„ | j        D ¦   «         ¦  «        f¦  «        S )Nc              3   ó4   K  — | ]}t          |¦  «        V — Œd S r   rs   rl   s     r   rn   z2NVMMADistributedLayout.__hash__.<locals>.<genexpr>6  ó(   è è € ÐAÐA¨#�5 ™:œ:ÐAÐAÐAÐAÐAÐAr   )ru   rt   r·   rP   r¸   rT   r   s    r   rv   zNVMMADistributedLayout.__hash__4  s_   € Ý•U˜4œ<Ñ(Ô(­%°Ô0BÑ*CÔ*CÅUÈ4ÔK[ÑE\ÔE\ÝÐAÐA°´ÐAÑAÔAÑAÔAðCñ Dô Dð 	Dr   c                 ó*   — t          | j        ¦  «        S r   )rX   rP   r   s    r   r   zNVMMADistributedLayout.rank8  s   € å�4Ô%Ñ&Ô&Ð&r   ry   r|   s   @r   r¶   r¶     sü   ø€ € € € € € ðð ð �#ŒYÐÐÑØ˜”9ÐÐÑØ�c”ÐÐÑØ"' %¸Ð"=Ñ"=Ô"=€J��T˜#”Y”Ð=Ð=Ñ=ðQð Qð Qð Qð Qð
ð 
ð 
ð^˜ð ^ð ^ð ^ð ^ðDð Dð Dð ð'ð 'ñ „Xð'ð 'ð 'ð 'ð 'r   r¶   c                   ó`   — e Zd ZdZed„ ¦   «         Zdee         defd„Z	dee         defd„Z
dS )ÚSharedLayoutz;
    Base class for shared memory layouts in Gluon IR.
    c                 ó    — t          | ¦  «        S r   r   r   s    r   r   zSharedLayout.typeB  r   r   r   r   c                 óB   — t          j        | d„ |D ¦   «         d¦  «        S )Nc                 ó,   — g | ]}t          |¦  «        ‘ŒS r   r   r   s     r   r"   z3SharedLayout.format_tensor_view.<locals>.<listcomp>G  r#   r   Fr$   r&   s     r   r'   zSharedLayout.format_tensor_viewF  r(   r   c                 óB   — t          j        | d„ |D ¦   «         d¦  «        S )Nc                 ó,   — g | ]}t          |¦  «        ‘ŒS r   r   r   s     r   r"   z5SharedLayout.format_hardware_view.<locals>.<listcomp>J  r#   r   Tr$   r&   s     r   r+   z!SharedLayout.format_hardware_viewI  r,   r   N)r-   r.   r/   r0   r1   r   r2   r3   r4   r'   r+   r   r   r   rÆ   rÆ   =  s“   € € € € € ðð ð ð$ð $ñ „Xð$ð_¨¨S¬	ð _°cð _ð _ð _ð _ð^¨$¨s¬)ð ^¸ð ^ð ^ð ^ð ^ð ^ð ^r   rÆ   c                 óú  — |s| S t          | ¦  «        }t          |d         ¦  «        }dg|z  }|D ]H}t          |¦  «        |k    sJ ‚t          |¦  «        D ]!}t          ||         ||         ¦  «        ||<   Œ"ŒIt          |¦  «        D ]$}||         dk    rd||<   Œ||xx         dz  cc<   Œ%t          |¦  «        D ]=}||         ||         z  dk    sJ d| › d|› �¦   «         ‚||xx         ||         z  cc<   Œ>|S )Nr   rŠ   r­   zShape z  is not divisible by CGA layout )r2   rX   ÚrangeÚmax)r   rT   Úshape_per_ctar   Ú	cga_shaperŽ   Úir   s           r   Ú_get_shape_per_ctarÒ   M  sS  € àð ØˆÝ˜‘K”K€MÝˆz˜!Œ}ÑÔ€DØ��d‘
€IØð 7ð 7ˆÝ�5‰zŒz˜TÒ!Ð!Ð!Ð!Ý�t‘”ð 	7ð 	7ˆAÝ˜y¨œ|¨U°1¬XÑ6Ô6ˆI�a‰LˆLð	7õ �4‰[Œ[ð ð ˆØ�QŒ<˜1ÒÐØˆI�a‰LˆLà�aˆLˆLŒL˜AÑˆLˆL‰LˆLÝ�T‰{Œ{ð .ð .ˆØ˜SÔ! I¨c¤NÑ2°aÒ7Ð7Ð7Ð9uÀ%Ð9uÐ9uÐisÐ9uÐ9uÑ7Ô7Ð7Ø�cÐÐÔ˜y¨œ~Ñ-ÐÐÑÐØÐr   c                   óÜ   ‡ — e Zd ZU dZeed<   eed<   dZeed<   dZeed<   dZ	eed<    e
e¬	¦  «        Zeee                  ed
<   ˆ fd„Zd„ Zeedd„¦   «         ¦   «         Zdefd„Zd„ Zˆ xZS )Ú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_widthÚelement_bitwidthr­   r   FÚ
transposedÚ
fp4_paddedrR   rT   c                 óV  •— t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         | j        pg }|r t          |d         ¦  «        | j	        k    sJ ‚t          ¦   «                              dt          | j	        ¦  «        ¦  «         t          ¦   «                              dt          |¦  «        ¦  «         | j        dv sJ ‚| j        d	v sJ ‚| j        r*| j        d
k    s
J d¦   «         ‚| j        dk    sJ d¦   «         ‚d S d S )NrÕ   rÖ   r×   rØ   r   r   rT   )é   é   é    é@   )r   rÜ   rÝ   é€   rÞ   z+fp4_padded only supports 128 byte swizzlingrÚ   z3fp4_padded is only supported for element_bitwidth=8)
rV   rW   r   rÕ   rÖ   r×   rØ   rT   rX   r   )r   rT   rY   s     €r   rZ   zNVMMASharedLayout.__post_init__x  s˜  ø€ Ý‰Œ×ÒÐ0Õ2FÀtÔG^Ñ2_Ô2_Ñ`Ô`Ð`Ý‰Œ×ÒÐ.Õ0DÀTÔEZÑ0[Ô0[Ñ\Ô\Ð\Ý‰Œ×Ò˜LÕ*>¸t¼Ñ*OÔ*OÑPÔPÐPÝ‰Œ×Ò˜LÕ*>¸t¼Ñ*OÔ*OÑPÔPÐPð ”_Ð*¨ˆ
Øð 	3Ý�z !”}Ñ%Ô%¨¬Ò2Ð2Ð2Ð2å‰Œ×Ò˜FÕ$8¸¼Ñ$CÔ$CÑDÔDÐDÝ‰Œ×Ò˜LÕ*>¸zÑ*JÔ*JÑKÔKÐKàÔ$¨Ð7Ð7Ð7Ð7ØÔ&Ð*:Ð:Ð:Ð:Ð:àŒ?ð 	eØÔ*¨cÒ1Ð1Ð1Ð3`Ñ1Ô1Ð1ØÔ(¨AÒ-Ð-Ð-Ð/dÑ-Ô-Ð-ð	eð 	eà-Ð-r   c                 ór   — |                      | j        | j        | j        | j        | j        | j        ¦  «        S r   )Úget_nvmma_shared_layoutrÕ   rÖ   r×   rØ   rT   r   r:   s     r   r<   zNVMMASharedLayout._to_ir�  s:   € Ø×.Ò.ØÔ#ØÔ!ØŒOØŒOØŒOØŒIñ
ô 
ð 	
r   Nc                 ó´  — |rdnd}|€| nt          | |¦  «        }t          | ¦  «        }|r|dd…         |dd…         z   }|d         |z  }||j        z  dz  }	|	dk    r|	dz  dk    rd}
n&|	dk    r|	dz  dk    rd}
n|	d	k    r|	d	z  dk    rd	}
nd}
d}|dd…         D ]}||z  }Œt          | ¦  «        dk     s|dk     rd}
t          |
|j        ||||¬
¦  «        S )zØReturns 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­   rŠ   NéÿÿÿÿrÚ   rÞ   r   rÝ   rÜ   )rÕ   rÖ   r   r×   rØ   rT   )rÒ   rX   Úprimitive_bitwidthrÔ   )Úblock_shapeÚdtyper×   rØ   rT   Úpacking_factorrÏ   r   Úcontig_dim_sizeÚcontig_dim_bytesrÕ   Úflatten_outer_dimÚsizes                r   Úget_default_forz!NVMMASharedLayout.get_default_for—  so  € ð )Ð/˜˜¨aˆØ'1Ð'9˜˜Õ?QÐR]Ð_iÑ?jÔ?jˆÝ�;ÑÔˆØð 	BØ)¨!¨"¨"Ô-°¸b¸q¸bÔ0AÑAˆMØ'¨Ô+¨nÑ<ˆØ*¨UÔ-EÑEÈÑJÐØ˜sÒ"Ð"Ð'7¸#Ñ'=ÀÒ'BÐ'BØ!$ÐÐØ Ò#Ð#Ð(8¸2Ñ(=ÀÒ(BÐ(BØ!#ÐÐØ Ò#Ð#Ð(8¸2Ñ(=ÀÒ(BÐ(BØ!#ÐÐà!"ÐàÐØ! # 2 #Ô&ð 	&ð 	&ˆDØ Ñ%ÐÐÝˆ{ÑÔ˜aÒÐÐ#4°qÒ#8Ð#8Ø!"Ðå Ø1Ø"Ô5ØØ!Ø!Ø!ð
ñ 
ô 
ð 	
r   r   c                 ó¨   — | j         r$d                     d„ | j         D ¦   «         ¦  «        nd}d| j        › d| j        › d| j        › d| j        › d|› d�S )Nra   c              3   óf   K  — | ],}d                       t          t          |¦  «        ¦  «        V — Œ-dS rj   rb   rl   s     r   rn   z+NVMMASharedLayout.mangle.<locals>.<genexpr>¿  ro   r   r`   ÚNVMMA_Ú_NVMMA)rT   rc   rÕ   rÖ   r×   rØ   rÀ   s     r   r?   zNVMMASharedLayout.mangle¾  s—   € ØUYÔUdÐl�S—X’XÐQÐQÀÄÐQÑQÔQÑQÔQÐQÐjlˆ
ð A˜Ô/ð  Að  A°$Ô2Gð  Að  AÈ$Ì/ð  Að  AÐ\`Ô\kð  Að  AÐnxð  Að  Að  Að  	Ar   c                 ó¨   — t          | j        | j        | j        | j        | j        | j        rt          d„ | j        D ¦   «         ¦  «        nd f¦  «        S )Nc              3   ó4   K  — | ]}t          |¦  «        V — Œd S r   rs   rl   s     r   rn   z-NVMMASharedLayout.__hash__.<locals>.<genexpr>Ä  rÃ   r   )ru   rÕ   rÖ   r   r×   rØ   rT   rt   r   s    r   rv   zNVMMASharedLayout.__hash__Â  s`   € Ý�TÔ,¨dÔ.CÀTÄYÐPTÔP_ÐaeÔapØEIÄ_Ð^•UÐAÐA°´ÐAÑAÔAÑAÔAÐAÐZ^ð`ñ aô að 	ar   )FFN)r-   r.   r/   r0   r3   rz   r   r×   ÚboolrØ   r   r2   rT   r   rZ   r<   Ústaticmethodr	   rë   r4   r?   rv   r{   r|   s   @r   rÔ   rÔ   d  s&  ø€ € € € € € ð
ð 
ð ÐÐÑØÐÐÑØ€Dˆ#€M€M�MØ€J�ÐÐÑØ€J�ÐÐÑØ"' %¸Ð"=Ñ"=Ô"=€J��T˜#”Y”Ð=Ð=Ñ=ðeð eð eð eð eð*
ð 
ð 
ð Øð#
ð #
ð #
ñ Ôñ „\ð#
ðJA˜ð Að Að Að Aðað að að að að að ar   rÔ   c                   óª   ‡ — e Zd ZU dZeed<   eed<   eed<   ee         ed<    ee¬¦  «        Z	eee                  ed<   ˆ fd„Z
d	„ Zd
efd„Zd„ Zˆ xZS )Ú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.
    rm   Ú	per_phaseÚ	max_phaserQ   rR   rT   c                 ó  •— t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         d S )Nrm   rö   r÷   rQ   rT   )rV   rW   r   rm   rö   r÷   rQ   rT   r‚   s    €r   rZ   z"SwizzledSharedLayout.__post_init__Ù  sÂ   ø€ Ý‰Œ×Ò˜EÕ#7¸¼Ñ#AÔ#AÑBÔBÐBÝ‰Œ×Ò˜KÕ)=¸d¼nÑ)MÔ)MÑNÔNÐNÝ‰Œ×Ò˜KÕ)=¸d¼nÑ)MÔ)MÑNÔNÐNÝ‰Œ×Ò˜GÕ%9¸$¼*Ñ%EÔ%EÑFÔFÐFÝ‰Œ×Ò˜LÕ*>¸t¼Ñ*OÔ*OÑPÔPÐPÐPÐPr   c                 óf   — |                      | j        | j        | j        | j        | j        ¦  «        S r   )Úget_swizzled_shared_layoutrm   rö   r÷   rQ   rT   r:   s     r   r<   zSwizzledSharedLayout._to_irà  s3   € Ø×1Ò1ØŒHØŒNØŒNØŒJØŒOñ
ô 
ð 	
r   r   c                 óÀ   — d„ }| j         r$d                     d„ | j         D ¦   «         ¦  «        nd}d| j        › d| j        › d| j        › d || j        ¦  «        › d|› d�S )Nc                 óZ   — | €dS d                      t          t          | ¦  «        ¦  «        S r_   rb   re   s    r   rg   z.SwizzledSharedLayout.mangle.<locals>.stringifyë  rh   r   ra   c              3   óf   K  — | ],}d                       t          t          |¦  «        ¦  «        V — Œ-dS rj   rb   rl   s     r   rn   z.SwizzledSharedLayout.mangle.<locals>.<genexpr>ð  ro   r   r`   ÚSSS_Ú_SSS)rT   rc   rm   rö   r÷   rQ   )r   rg   rT   s      r   r?   zSwizzledSharedLayout.mangleé  s‰   € ð	)ð 	)ð 	)ð
 VZÔUdÐl�S—X’XÐQÐQÀÄÐQÑQÔQÑQÔQÐQÐjlˆ
Øk�d”hÐkÐk ¤ÐkÐk°$´.ÐkÐkÀ9À9ÈTÌZÑCXÔCXÐkÐkÐ[eÐkÐkÐkÐkr   c           
      ó¤   — t          | j        | j        | j        t	          | j        ¦  «        t	          d„ | j        D ¦   «         ¦  «        f¦  «        S )Nc              3   ó4   K  — | ]}t          |¦  «        V — Œd S r   rs   rl   s     r   rn   z0SwizzledSharedLayout.__hash__.<locals>.<genexpr>õ  s-   è è € ÐOvÐOvÐ_bÕPUÐVYÑPZÔPZÐOvÐOvÐOvÐOvÐOvÐOvr   )ru   rm   rö   r÷   rt   rQ   rT   r   s    r   rv   zSwizzledSharedLayout.__hash__ó  sP   € ÝØŒX�t”~ t¤~µu¸T¼ZÑ7HÔ7HÍ%ÐOvÐOvÐfjÔfuÐOvÑOvÔOvÑJvÔJvÐwñyô yð 	yr   )r-   r.   r/   r0   r3   rz   r   r   r2   rT   rZ   r<   r4   r?   rv   r{   r|   s   @r   rõ   rõ   Ç  sà   ø€ € € € € € ð	ð 	ð 
€H€H�HØ€N€N�NØ€N€N�NØ�Œ9ÐÐÑØ"' %¸Ð"=Ñ"=Ô"=€J��T˜#”Y”Ð=Ð=Ñ=ðQð Qð Qð Qð Qð
ð 
ð 
ðl˜ð lð lð lð lðyð yð yð yð yð yð yr   rõ   c                   óè   ‡ — e Zd ZU dZeee                  ed<   eee                  ed<   eee                  ed<   ee         ed<   ˆ fd„Zd„ Zde	fd	„Z
d
„ Zeeg fd„¦   «         ¦   «         Zd„ Zˆ xZS )Ú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_basesrT   r   c                 ót  •— t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          | j        ¦  «        }| j        D ]}t          |¦  «        |k    sJ ‚Œ| j        D ]}t          |¦  «        |k    sJ ‚Œ|                      ¦   «          d S )Nr  r  rT   r   )	rV   rW   r   r  r  rT   r   rX   Úverifyr˜   s      €r   rZ   z PaddedSharedLayout.__post_init__8  s  ø€ Ý‰Œ×ÒÐ4µmÀDÔD_Ñ6`Ô6`ÑaÔaÐaÝ‰Œ×Ò˜N­M¸$Ô:KÑ,LÔ,LÑMÔMÐMÝ‰Œ×Ò˜L­-¸¼Ñ*HÔ*HÑIÔIÐIÝ‰Œ×Ò˜G¥]°4´:Ñ%>Ô%>Ñ?Ô?Ð?å�4”:‰ŒˆàÔ&ð 	&ð 	&ˆEÝ�u‘:”: Ò%Ð%Ð%Ð%Ð%Ø”_ð 	&ð 	&ˆEÝ�u‘:”: Ò%Ð%Ð%Ð%Ð%à�Š‰Œˆˆˆr   c                 ót   — t          | j        Ž \  }}|                     ||| j        | j        | j        ¦  «        S r   )Úzipr  Úget_padded_shared_layoutr  rT   r   )r   r;   Ú	intervalsÚpaddingss       r   r<   zPaddedSharedLayout._to_irG  s;   € Ý! 4Ô#>Ð?Ñˆ	�8Ø×/Ò/°	¸8ÀTÔEVÐX\ÔXgÐimÔisÑtÔtÐtr   r   c           	      óH   — d| j         › d| j        › d| j        › d| j        › d�	S )NÚPaddedShared_ra   Ú_PaddedShared)r  r  rT   r   r   s    r   r?   zPaddedSharedLayout.mangleK  s9   € Ø|˜tÔ:Ð|Ð|¸TÔ=NÐ|Ð|ÐQUÔQ`Ð|Ð|ÐcgÔcmÐ|Ð|Ð|Ð|r   c                 óü  ‡— | j         }t          |¦  «        dk    s
J d¦   «         ‚t          d„ |D ¦   «         ¦  «        sJ ‚t          |Ž \  }}t	          t          |¦  «        ¦  «        }t          |¦  «        t          |¦  «        k    sJ ‚d„ Št          ˆfd„|D ¦   «         ¦  «        s
J d¦   «         ‚t          ˆfd„|D ¦   «         ¦  «        s
J d¦   «         ‚t          | j        ¦  «        }|dk    s
J d	¦   «         ‚d S )
Nr   zVPaddedSharedLayout interval_padding_pairs must have at least one interval-padding pairc              3   ó<   K  — | ]}t          |¦  «        d k    V — ŒdS )r­   Nr¬   )r    Úpairs     r   rn   z,PaddedSharedLayout.verify.<locals>.<genexpr>Q  s,   è è € Ð4Ð4 d•3�t‘9”9 ’>Ð4Ð4Ð4Ð4Ð4Ð4r   c                 ó&   — | dk    o| | dz
  z  dk    S ©Nr   rŠ   r   ©Úns    r   ú<lambda>z+PaddedSharedLayout.verify.<locals>.<lambda>W  ó   €  ! a¢%Ð"<¨A°°Q±©K¸1Ò,<€ r   c              3   ó.   •K  — | ]} ‰|¦  «        V — Œd S r   r   ©r    r  Úis_power_of_2s     €r   rn   z,PaddedSharedLayout.verify.<locals>.<genexpr>X  s-   øè è € Ð7Ð7¨�=�= Ñ#Ô#Ð7Ð7Ð7Ð7Ð7Ð7r   z;PaddedSharedLayout interval values must all be power of twoc              3   ó.   •K  — | ]} ‰|¦  «        V — Œd S r   r   r  s     €r   rn   z,PaddedSharedLayout.verify.<locals>.<genexpr>Y  s-   øè è € Ð6Ð6¨�=�= Ñ#Ô#Ð6Ð6Ð6Ð6Ð6Ð6r   z:PaddedSharedLayout padding values must all be power of twoz*PaddedSharedLayout order must not be empty)r  rX   r¯   r	  r2   Úsetr   )r   Úpairsr  r  Úunique_intervalsr   r  s         @r   r  zPaddedSharedLayout.verifyN  s  ø€ ØÔ+ˆÝ�5‰zŒz˜AŠ~ˆ~ˆ~Ðw‰~Œ~ˆ~ÝÐ4Ð4¨eÐ4Ñ4Ô4Ñ4Ô4Ð4Ð4Ð4Ý! 5˜kÑˆ	�8å¥ I¡¤Ñ/Ô/ÐÝÐ#Ñ$Ô$­¨I©¬Ò6Ð6Ð6Ð6à<Ð<ˆÝÐ7Ð7Ð7Ð7¨YÐ7Ñ7Ô7Ñ7Ô7ÐvÐvÐ9vÑvÔvÐ7ÝÐ6Ð6Ð6Ð6¨XÐ6Ñ6Ô6Ñ6Ô6ÐtÐtÐ8tÑtÔtÐ6å�4”:‰ŒˆØ�aŠxˆxˆxÐE‰xŒxˆxˆxˆxr   c           	      óÐ  ‡‡‡	— t          |¦  «        t          |¦  «        k    sJ ‚d„ Š	t          ˆ	fd„|D ¦   «         ¦  «        sJ ‚t          |¦  «        }|rt          ||¦  «        n|}g }|D ]hŠt          t	          t          j        |‰         ¦  «        ¦  «        ¦  «        D ]1Š|                     ˆˆfd„t          |¦  «        D ¦   «         ¦  «         Œ2Œit          | |||¦  «        S )zžReturns a PaddedSharedLayout with the given interval and padding pairs and an identity mapping as the linear component for the given shape and order.
        c                 ó&   — | dk    o| | dz
  z  dk    S r  r   r  s    r   r  z6PaddedSharedLayout.with_identity_for.<locals>.<lambda>d  r  r   c              3   ó.   •K  — | ]} ‰|¦  «        V — Œd S r   r   r  s     €r   rn   z7PaddedSharedLayout.with_identity_for.<locals>.<genexpr>e  s-   øè è € Ð3Ð3¨�=�= Ñ#Ô#Ð3Ð3Ð3Ð3Ð3Ð3r   c                 ó*   •— g | ]}|‰k    rd ‰z  nd‘ŒS )rŠ   r   r   )r    rÑ   rŽ   r   s     €€r   r"   z8PaddedSharedLayout.with_identity_for.<locals>.<listcomp>m  s)   ø€ Ð$XÐ$XÐ$XÀq°1¸²8°8 Q¨%¡Z ZÀÐ$XÐ$XÐ$Xr   )	rX   r¯   rÒ   rÍ   r3   ÚmathÚlog2r°   r  )
r  r   rQ   rT   r   rÏ   r  rŽ   r   r  s
          @@@r   Úwith_identity_forz$PaddedSharedLayout.with_identity_for^  s  øøø€ õ
 �5‰zŒz�S ™ZœZÒ'Ð'Ð'Ð'Ø<Ð<ˆÝÐ3Ð3Ð3Ð3¨UÐ3Ñ3Ô3Ñ3Ô3Ð3Ð3Ð3å�5‰zŒzˆØAKÐVÕ*¨5°*Ñ=Ô=Ð=ÐQVˆàˆØð 	Zð 	ZˆCÝ�s¥4¤9¨]¸3Ô-?Ñ#@Ô#@ÑAÔAÑBÔBð Zð Z�Ø×#Ò#Ð$XÐ$XÐ$XÐ$XÐ$XÍEÐRVÉKÌKÐ$XÑ$XÔ$XÑYÔYÐYÐYðZõ "Ð"8¸,È
ÐTYÑZÔZÐZr   c           
      ó*  — t          t          t          t          | j        ¦  «        ¦  «        t          t          t          | j        ¦  «        ¦  «        t          t          t          | j        ¦  «        ¦  «        t          | j        ¦  «        f¦  «        S r   )ru   rt   rd   r  r  rT   r   r   s    r   rv   zPaddedSharedLayout.__hash__q  so   € Ý•U�3�u dÔ&AÑBÔBÑCÔCÅUÍ3ÍuÐVZÔVgÑKhÔKhÑEiÔEiÝ�3�u d¤oÑ6Ô6Ñ7Ô7½¸t¼zÑ9JÔ9JðLñ Mô Mð 	Mr   )r-   r.   r/   r0   r   r3   rz   rZ   r<   r4   r?   r  ró   r	   r&  rv   r{   r|   s   @r   r  r  ø  s  ø€ € € € € € ð8ð 8ðr !  c¤œOÐ+Ð+Ñ+Ø�t˜C”y”/Ð!Ð!Ñ!Ø�T˜#”Y”ÐÐÑØ�Œ9ÐÐÑðð ð ð ð ðuð uð uð}˜ð }ð }ð }ð }ðFð Fð Fð  ØØKMð [ð [ð [ñ Ôñ „\ð[ð"Mð Mð Mð Mð Mð Mð Mr   r  c                   ó¼   ‡ — e Zd ZU dZeee                  ed<    ee¬¦  «        Z	eee                  ed<   dZ
eed<   ˆ fd„Zd„ Zd	efd
„Zed„ ¦   «         Zd„ Zˆ xZS )ÚSharedLinearLayoutzGRepresents a shared memory layout defined via an explicit LinearLayout.r  rR   r–   rÛ   Ú	alignmentc                 óª  •— t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          ¦   «                              dt          | j        ¦  «        ¦  «         t          | j        ¦  «        dk    s
J d¦   «         ‚t          | j        d         ¦  «        }|dk    s
J d¦   «         ‚| j        D ]}t          |¦  «        |k    sJ ‚Œ| j        D ]}t          |¦  «        |k    sJ ‚Œ| j        dk    r| j        | j        dz
  z  dk    s
J d¦   «         ‚d S )Nr  r–   r*  r   z1SharedLinearLayout offset_bases must not be emptyrŠ   z<SharedLinearLayout alignment must be a positive power of two)rV   rW   r   r  r–   r   r*  rX   r˜   s      €r   rZ   z SharedLinearLayout.__post_init__~  sV  ø€ Ý‰Œ×Ò˜N­M¸$Ô:KÑ,LÔ,LÑMÔMÐMÝ‰Œ×Ò˜M­=¸Ô9IÑ+JÔ+JÑKÔKÐKÝ‰Œ×Ò˜KÕ)=¸d¼nÑ)MÔ)MÑNÔNÐNå�4Ô$Ñ%Ô%¨Ò*Ð*Ð*Ð,_Ñ*Ô*Ð*Ý�4Ô$ QÔ'Ñ(Ô(ˆØ�aŠxˆxˆxÐL‰xŒxˆxØÔ&ð 	&ð 	&ˆEÝ�u‘:”: Ò%Ð%Ð%Ð%Ð%ØÔ%ð 	&ð 	&ˆEÝ�u‘:”: Ò%Ð%Ð%Ð%Ð%ØŒ~ Ò!Ð! t¤~¸¼È!Ñ9KÑ'LÐQRÒ&RÐ&RÐ&RØJñ 'SÔ&RÐRÐ&RÐ&Rr   c                 óN   — |                      | j        | j        | j        ¦  «        S r   )Úget_shared_linear_layoutr  r–   r*  r:   s     r   r<   zSharedLinearLayout._to_ir�  s$   € Ø×/Ò/°Ô0AÀ4ÔCSÐUYÔUcÑdÔdÐdr   r   c                 ó8   — d| j         › d| j        › d| j        › d�S )NÚSharedLinear_ra   Ú_SharedLinear)r  r–   r*  r   s    r   r?   zSharedLinearLayout.mangle�  s+   € Øc˜tÔ0ÐcÐc°4Ô3CÐcÐcÀdÄnÐcÐcÐcÐcr   c                 óú   — t          | j        d         ¦  «        }dg|z  }t          j        | j        | j        ¦  «        D ]0}t          |¦  «        D ]\  }}t          ||         |¦  «        ||<   ŒŒ1d„ |D ¦   «         S )Nr   rŠ   c                 ó   — g | ]}d |z  ‘ŒS )r­   r   r   s     r   r"   z,SharedLinearLayout.shape.<locals>.<listcomp>š  s   € Ð*Ð*Ð*˜!��A‘Ð*Ð*Ð*r   )rX   r  Ú	itertoolsÚchainr–   Ú	enumeraterÎ   )r   r   Ú
max_strideÚbrÑ   Úbis         r   r   zSharedLinearLayout.shape“  sŽ   € å�4Ô$ QÔ'Ñ(Ô(ˆØ�S˜4‘Zˆ
Ý” Ô!2°DÔ4DÑEÔEð 	7ð 	7ˆAÝ" 1™œð 7ð 7‘��2Ý # J¨q¤M°2Ñ 6Ô 6�
˜1‘�ð7à*Ð*˜zÐ*Ñ*Ô*Ð*r   c           	      óÄ   — t          t          t          t          | j        ¦  «        ¦  «        t          t          t          | j        ¦  «        ¦  «        | j        f¦  «        S r   )ru   rt   rd   r  r–   r*  r   s    r   rv   zSharedLinearLayout.__hash__œ  sN   € ÝÝ•#•e˜TÔ.Ñ/Ô/Ñ0Ô0Ý•#•e˜TÔ-Ñ.Ô.Ñ/Ô/ØŒNð
ñ ô ð 	r   )r-   r.   r/   r0   r   r3   rz   r   r2   r–   r*  rZ   r<   r4   r?   r1   r   rv   r{   r|   s   @r   r)  r)  v  sê   ø€ € € € € € àQÐQà�t˜C”y”/Ð!Ð!Ñ!Ø#( 5¸Ð#>Ñ#>Ô#>€K��d˜3”i”Ð>Ð>Ñ>Ø€IˆsÐÐÑðKð Kð Kð Kð Kðeð eð eðd˜ð dð dð dð dð ð+ð +ñ „Xð+ðð ð ð ð ð ð r   r)  c                 óÌ   — dg|z  }| s|S d }| D ]T}t          d„ t          |¦  «        D ¦   «         d ¦  «        }|�|}||xx         dz  cc<   Œ>|s|€J ‚||xx         dz  cc<   ŒU|S )NrŠ   c              3   ó,   K  — | ]\  }}|d k    ¯|V — ŒdS )r   Nr   )r    rÑ   Úvs      r   rn   z bases_per_dim.<locals>.<genexpr>¯  s*   è è € Ð=Ð=™$˜!˜Q°a¸1²f°f�A°f°f°f°fÐ=Ð=r   r­   )Únextr5  )Úbasesr   Úskip_broadcastÚresultÚnon_zero_idxrŽ   Úidxs          r   Úbases_per_dimrC  ¥  s¬   € ØˆS�4‰Z€Fàð Øˆà€Làð 	&ð 	&ˆåÐ=Ð=¥)¨EÑ"2Ô"2Ð=Ñ=Ô=¸tÑDÔDˆØˆ?ØˆLØ�3ˆKˆKŒK˜1ÑˆKˆK‰KˆKØð 	&àÐ+Ð+Ð+Ø�<Ð Ð Ô  AÑ%Ð Ð Ñ øà€Mr   c                 óà   — t          | t          ¦  «        r"t          | j        t	          |¦  «        ¦  «        S t          | t
          t          f¦  «        rt          | j        |¦  «        S | j        S r   )	Ú
isinstancer’   rC  r•   rX   r~   r    rP   r€   )Úlayoutr   s     r   rP   rP   »  sa   € Ý�&Õ1Ñ2Ô2ð $Ý˜VÔ.µ°E±
´
Ñ;Ô;Ð;Ý	�F�[Õ*:Ð;Ñ	<Ô	<ð $Ý˜Vœ]¨EÑ2Ô2Ð2àÔ#Ð#r   )T)Údataclassesr   r   r3  r$  Útypingr   Útriton.language.corer   r   r   Útriton.runtime.jitr	   Útriton._C.libtritonr
   r   r7   rF   rM   r~   r’   r    r¶   rÆ   rÒ   rÔ   rõ   r  r)  rC  rP   r   r   r   ú<module>rL     sa  ðØ (Ð (Ð (Ð (Ð (Ð (Ð (Ð (Ø Ð Ð Ð Ø €€€Ø Ð Ð Ð Ð Ð à TÐ TÐ TÐ TÐ TÐ TÐ TÐ TÐ TÐ TØ 1Ð 1Ð 1Ð 1Ð 1Ð 1Ø (Ð (Ð (Ð (Ð (Ð (ð^ð ^ð ^ð ^ð ^ñ ^ô ^ð ^ð( €�$ÐÑÔð
3ð 
3ð 
3ð 
3ð 
3Ð"ñ 
3ô 
3ñ Ôð
3ð €�$ÐÑÔð
8ð 
8ð 
8ð 
8ð 
8Ð'ñ 
8ô 
8ñ Ôð
8ð €�$ÐÑÔð:ð :ð :ð :ð :Ð%ñ :ô :ñ Ôð:ðz €�$ÐÑÔð'Xð 'Xð 'Xð 'Xð 'XÐ#ñ 'Xô 'Xñ Ôð'XðT €�$ÐÑÔð6ð 6ð 6ð 6ð 6Ð/ñ 6ô 6ñ Ôð6ðr €�$ÐÑÔð0ð 0ð 0ð 0ð 0Ð(ñ 0ô 0ñ Ôð0ðf €�$˜4Ð Ñ Ô ð''ð ''ð ''ð ''ð ''Ð.ñ ''ô ''ñ !Ô ð''ðT^ð ^ð ^ð ^ð ^ñ ^ô ^ð ^ð  ðð ñ Ôðð, €�$ÐÑÔð_að _að _að _að _a˜ñ _aô _añ Ôð_aðD €�$˜4Ð Ñ Ô ð-yð -yð -yð -yð -y˜<ñ -yô -yñ !Ô ð-yð` €�$˜4Ð Ñ Ô ðzMð zMð zMð zMð zM˜ñ zMô zMñ !Ô ðzMðz €�$ÐÑÔð*ð *ð *ð *ð *˜ñ *ô *ñ Ôð*ð\ð ð ð ð,$ð $ð $ð $ð $r   