Ë
    ŒÓ^jc  ã                  ó¬   — d dl mZ d dlmZmZ d dlmZmZ d dlm	Z	 d dl
mZ ddgZ ed¬	«       G d
„ de«      «       Z ed¬	«       G d„ de«      «       Zy)é    )Úannotations)Ú	dataclassÚfield)ÚListÚOptional)Ú_unwrap_if_constexpr)ÚDistributedLayoutÚAMDMFMALayoutÚAMDWMMALayoutT)Úfrozenc                  ó®   ‡ — e Zd ZU dZded<   ded<   ded<   ded<   d	Zd
ed<   d	Zded<    ee¬«      Z	ded<   ˆ fd„Z
d„ Zdd„Zd„ Zd„ Zed„ «       Zˆ xZS )r
   a’  
    Represents a layout for AMD MFMA (matrix core) operations.

    Args:
        version (int): The GPU architecture.
        instr_shape (List[int]): The shape in the form of (M, N, K) of the matrix.
        transposed (bool): Indicates the result tensor is transposed so that each thread holds consecutive elements in the same row instead of column, which is good for chained dot and global write.
        warps_per_cta (List[int]): The warp layout in the block.
        element_bitwidth Optional(int): Bit width of the output element type. Supported values are 32 and 64. Defaults to 32.
        tiles_per_warp Optional(List[int]): The tile layout within a warp. Defaults to unit tile layout, i.e., single tile on all dimensions.
        cga_layout (Optional[List[List[int]]]): Bases describing CTA tiling.

    Current supported versions:

    - 1: gfx908
    - 2: gfx90a
    - 3: gfx942
    - 4: gfx950
    ÚintÚversionz	List[int]Úinstr_shapeÚboolÚ
transposedÚwarps_per_ctaNúOptional[int]Úelement_bitwidthúOptional[List[int]]Útiles_per_warp©Údefault_factoryúList[List[int]]Ú
cga_layoutc                ó®  •— t         ‰| �  dt        | j                  «      «       t         ‰| �  dt        | j                  «      «       t         ‰| �  dt        | j
                  «      «       t         ‰| �  dt        | j                  «      «       t         ‰| �  dt        | j                  «      «       t         ‰| �  dt        | j                  «      «       t         ‰| �  dt        | j                  «      «       | j                  €t         ‰| �  dd«       | j                  €'t         ‰| �  dd	gt        | j                  «      z  «       | j                  «        y )
Nr   r   r   r   r   r   r   é    é   )ÚsuperÚ__setattr__r   r   r   r   r   r   r   r   ÚlenÚverify)ÚselfÚ	__class__s    €úz/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/amd/_layouts.pyÚ__post_init__zAMDMFMALayout.__post_init__,   s  ø€ Ü‰Ñ˜IÔ';¸D¿L¹LÓ'IÔJÜ‰Ñ˜MÔ+?À×@PÑ@PÓ+QÔRÜ‰Ñ˜LÔ*>¸t¿¹Ó*OÔPÜ‰Ñ˜OÔ-AÀ$×BTÑBTÓ-UÔVÜ‰ÑÐ.Ô0DÀT×EZÑEZÓ0[Ô\Ü‰ÑÐ,Ô.BÀ4×CVÑCVÓ.WÔXÜ‰Ñ˜LÔ*>¸t¿¹Ó*OÔPà× Ñ Ð(Ü‰GÑÐ 2°BÔ7Ø×ÑÐ&Ü‰GÑÐ 0°1°#¼¸D×<NÑ<NÓ8OÑ2OÔPà�‰�ó    c           	     ó¼   — |j                  | j                  | j                  | j                  | j                  | j
                  | j                  | j                  «      S ©N)Úget_amd_mfma_layoutr   r   r   r   r   r   r   ©r#   Úbuilders     r%   Ú_to_irzAMDMFMALayout._to_ir<   sN   € Ø×*Ñ*Ø�L‰LØ×ÑØ×ÑØ�O‰OØ�O‰OØ×ÑØ×!Ñ!ó
ð 	
r'   c                ót  — d„ } || j                   r7| j                   D �cg c]!  }dj                  t        t        |«      «      ‘Œ# c}nd «      }d| j                  › d || j
                  «      › d| j                  › d || j                  «      › d| j                  › d || j                  «      › d|› d�S c c}w )Nc                óF   — | €ydj                  t        t        | «      «      S ©NÚ Ú_©ÚjoinÚmapÚstr©Úxs    r%   Ú	stringifyz'AMDMFMALayout.mangle.<locals>.stringifyI   ó   € ØˆyØØ—8‘8œC¤ Q›KÓ(Ð(r'   ú~ÚMFMA_r2   Ú_MFMA)
r   r4   r5   r6   r   r   r   r   r   r   )r#   r9   Úvecr   s       r%   ÚmanglezAMDMFMALayout.mangleG   sô   € ò	)ñ
 ÐW[×WfÒWfÀ4Ç?Á?ÖS¸C §¡¬¬S°#«Õ 7ÓSÐlpÓqˆ
Ø�t—|‘|�n A¡i°×0@Ñ0@Ó&AÐ%BÀ!ÀDÇOÁOÐCTÐTUÑV_Ð`d×`rÑ`rÓVsÐUtÐtuÐvz÷  wLñ  wLð  vMð  MNñ  OXð  Y]÷  Ylñ  Yló  Omð  Nnð  noð  pzð  o{ð  {@ð  Að  	Aùò  Ts    &B5c                ó’  ‡— | j                   dk\  r| j                   dk  sJ d«       ‚t        | j                  «      dk(  sJ d«       ‚ddgddgddgddgg}| j                  d	d
 |v sJ d| j                  › �«       ‚| j                  dv sJ d«       ‚t        | j                  «      Št        ˆfd„| j                  D «       «      sJ d«       ‚y )Nr   é   z#version must be in the [1, 4] rangeé   z,instr_shape must follow the (M, N, K) formatr   é   é@   r   é   zinvalid intrinsic shape )r   rD   z!element bitwidth must be 32 or 64c              3  ó:   •K  — | ]  }t        |«      ‰k(  –— Œ y ­wr)   )r!   )Ú.0r>   Úranks     €r%   ú	<genexpr>z'AMDMFMALayout.verify.<locals>.<genexpr>Y   s   øè ø€ Ò?¨”3�s“8˜tÕ#Ñ?ùs   ƒúcga_layout basis rank mismatch)r   r!   r   r   r   Úallr   )r#   Úvalid_shapesrH   s     @r%   r"   zAMDMFMALayout.verifyQ   sÛ   ø€ Ø�|‰|˜qÒ  T§\¡\°QÒ%6Ð]Ð8]Ó]Ð6Ü�4×#Ñ#Ó$¨Ò)ÐYÐ+YÓYÐ)Ø˜R˜ 2 r (¨R°¨G°a¸°WÐ=ˆØ×Ñ  !Ð$¨Ñ4ÐcÐ8PÐQU×QaÑQaÐPbÐ6cÓcÐ4Ø×$Ñ$¨Ñ0ÐUÐ2UÓUÐ0ä�4×%Ñ%Ó&ˆÜÓ?¨t¯©Ô?Ô?ÐaÐAaÓaÑ?r'   c                ó@  — t        | j                  t        | j                  «      | j                  t        | j
                  «      | j                  r| j                  nd | j                  rt        | j                  «      nd t        d„ | j                  D «       «      f«      S )Nc              3  ó2   K  — | ]  }t        |«      –— Œ y ­wr)   ©Útuple©rG   r>   s     r%   rI   z)AMDMFMALayout.__hash__.<locals>.<genexpr>c   ó   è ø€ Ò8 ”%˜—*Ñ8ùó   ‚)	Úhashr   rP   r   r   r   r   r   r   ©r#   s    r%   Ú__hash__zAMDMFMALayout.__hash__[   s{   € ÜØ�L‰LÜ�$×"Ñ"Ó#Ø�O‰OÜ�$×$Ñ$Ó%Ø%)×%:Ò%:ˆD×!Ò!ÀØ*.×*=Ò*=ŒE�$×%Ñ%Ô&À4ÜÑ8¨¯©Ô8Ó8ð
ó ð 	r'   c                ó,   — t        | j                  «      S r)   )r!   r   rU   s    r%   rH   zAMDMFMALayout.rankf   s   € ä�4×%Ñ%Ó&Ð&r'   ©Úreturnr6   )Ú__name__Ú
__module__Ú__qualname__Ú__doc__Ú__annotations__r   r   r   Úlistr   r&   r-   r?   r"   rV   ÚpropertyrH   Ú__classcell__©r$   s   @r%   r
   r
      su   ø… ñð& ƒLØÓØÓØÓØ&*Ð�mÓ*Ø*.€NÐ'Ó.Ù"'¸Ô"=€J�Ó=ôò 	
óAòbò	ð ñ'ó ô'r'   c                  ó¢   ‡ — e Zd ZU dZded<   ded<   ded<   dZd	ed
<   dZded<    ee¬«      Z	ded<   dZ
ded<   ˆ fd„Zd„ Zdd„Zd„ Zd„ Zˆ xZS )r   að  
    Represents a layout for AMD WMMA (matrix core) operations.

    Args:
        version (int): Indicates the GPU architecture.
        transposed (bool): Indicates the result tensor is transposed.
        warp_bases (List[List[int]]): Warp bases for CTA layout.
        reg_bases (Optional[List[List[int]]]): Repetition (register) bases for CTA layout.
        instr_shape (Optional[List[int]]): Instruction shape (M, N, K). Defaults to (16, 16, 16).
        cga_layout (Optional[List[List[int]]]): Bases describing CTA tiling.
        rank (Optional[int]): rank of warp and register bases. Default to 2 if missing.

    Current supported versions:

    - 1: RDNA3; e.g., gfx1100, gfx1101
    - 2: RDNA4; e.g., gfx1200, gfx1201
    - 3: gfx1250
    r   r   r   r   r   Ú
warp_basesNzOptional[List[List[int]]]Ú	reg_basesr   r   r   r   r   rH   c           	     óò  •— t         ‰| �  dt        | j                  «      «       t         ‰| �  dt        | j                  «      «       t         ‰| �  dt        | j
                  «      D �cg c]  }t        |«      ‘Œ c}«       t         ‰| �  d| j                  �,t        | j                  «      D �cg c]  }t        |«      ‘Œ c}ng «       | j                  �t        | j                  «      ng d¢}t         ‰| �  dt        |«      «       t         ‰| �  dt        | j                  «      «       | j                  �t        | j                  «      nd}t         ‰| �  d	|«       | j                  «        y c c}w c c}w )
Nr   r   rd   re   )rC   rC   rC   r   r   rE   rH   )r   r    r   r   r   rd   r_   re   r   r   rH   r"   )r#   Úinnerr   rH   r$   s       €r%   r&   zAMDWMMALayout.__post_init__‡   s3  ø€ Ü‰Ñ˜IÔ';¸D¿L¹LÓ'IÔJÜ‰Ñ˜LÔ*>¸t¿¹Ó*OÔPÜ‰Ñ˜LÔDXÐY]×YhÑYhÓDiÖ*j¸5¬4°­;Ò*jÔkÜ‰Ñ˜KàSW×SaÑSaÐSmÔ*>¸t¿~¹~Ó*NöPØ!&ô " %�[ó PØsuô	wð AE×@PÑ@PÐ@\Ô*¨4×+;Ñ+;Ô<ÒbnˆÜ‰Ñ˜MÔ+?ÀÓ+LÔMÜ‰Ñ˜LÔ*>¸t¿¹Ó*OÔPØ26·)±)Ð2GÔ# D§I¡IÔ.ÈQˆÜ‰Ñ˜F DÔ)Ø�‰�ùò +kùòPs   Á)E/
Â/E4
c           	     ó¼   — |j                  | j                  | j                  | j                  | j                  | j
                  | j                  | j                  «      S r)   )Úget_amd_wmma_layoutr   r   rd   re   r   r   rH   r+   s     r%   r-   zAMDWMMALayout._to_ir•   sH   € Ø×*Ñ*Ø�L‰LØ�O‰OØ�O‰OØ�N‰NØ�O‰OØ×ÑØ�I‰Ió
ð 	
r'   c                ó  ‡— d„ Šˆfd„} || j                   «      } || j                  «      } || j                  «      }d| j                  › d| j                  › d|› d|› d ‰| j
                  «      › d|› d| j                  › d�S )Nc                óF   — | €ydj                  t        t        | «      «      S r0   r3   r7   s    r%   r9   z'AMDWMMALayout.mangle.<locals>.stringify¢   r:   r'   c                ó„   •—  ‰| r1| D �cg c]!  }dj                  t        t        |«      «      ‘Œ# c}«      S d «      S c c}w )Nr;   r3   )r8   r>   r9   s     €r%   Únested_stringifyz.AMDWMMALayout.mangle.<locals>.nested_stringify§   s5   ø€ ÙÁqÀÖC¸#˜cŸh™h¤s¬3°£}Õ5ÒCÓSÐSÈdÓSÐSùÒCs   Š&=ÚWMMA_r2   Ú_WMMA)rd   re   r   r   r   r   rH   )r#   rm   rd   re   r   r9   s        @r%   r?   zAMDWMMALayout.mangle    s«   ø€ ò	)ô
	Tñ & d§o¡oÓ6ˆ
Ù$ T§^¡^Ó4ˆ	Ù% d§o¡oÓ6ˆ
Ø�t—|‘|�n A d§o¡oÐ%6°a¸
°|À1ÀYÀKÈqÑQZÐ[_×[kÑ[kÓQlÐPmÐmnÐoyÐnzÐz{ð  }A÷  }Fñ  }Fð  |Gð  GLð  Mð  	Mr'   c                ó"  ‡ — ‰ j                   dk\  r‰ j                   dk  sJ d«       ‚t        ‰ j                  «      dkD  r,t        ‰ j                  d   «      ‰ j                  k(  sJ d«       ‚t	        ˆ fd„‰ j
                  D «       «      sJ d«       ‚y )Nr   rB   z#version must be in the [1, 3] ranger   zwarp_bases basis rank mismatchc              3  óN   •K  — | ]  }t        |«      ‰j                  k(  –— Œ y ­wr)   )r!   rH   )rG   r>   r#   s     €r%   rI   z'AMDWMMALayout.verify.<locals>.<genexpr>³   s   øè ø€ ÒD¨S”3�s“8˜tŸy™yÕ(ÑDùs   ƒ"%rJ   )r   r!   rd   rH   rK   r   rU   s   `r%   r"   zAMDWMMALayout.verify¯   s{   ø€ Ø�|‰|˜qÒ  T§\¡\°QÒ%6Ð]Ð8]Ó]Ð6Üˆt�‰Ó !Ò#Ü�t—‘ qÑ)Ó*¨d¯i©iÒ7ÐYÐ9YÓYÐ7ÜÓD°D·O±OÔDÔDÐfÐFfÓfÑDr'   c                ó@  — t        | j                  | j                  t        d„ | j                  D «       «      t        d„ | j
                  D «       «      | j                  rt        | j                  «      nd t        d„ | j                  D «       «      | j                  f«      S )Nc              3  ó2   K  — | ]  }t        |«      –— Œ y ­wr)   rO   rQ   s     r%   rI   z)AMDWMMALayout.__hash__.<locals>.<genexpr>¹   rR   rS   c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wr)   rO   rQ   s     r%   rI   z)AMDWMMALayout.__hash__.<locals>.<genexpr>º   s   è ø€ Ò7 ”%˜—*Ñ7ùrS   c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wr)   rO   rQ   s     r%   rI   z)AMDWMMALayout.__hash__.<locals>.<genexpr>¼   rR   rS   )	rT   r   r   rP   rd   re   r   r   rH   rU   s    r%   rV   zAMDWMMALayout.__hash__µ   sw   € ÜØ�L‰LØ�O‰OÜÑ8¨¯©Ô8Ó8ÜÑ7¨¯©Ô7Ó7Ø'+×'7Ò'7ŒE�$×"Ñ"Ô#¸TÜÑ8¨¯©Ô8Ó8Ø�I‰Ið
ó ð 	r'   rX   )rZ   r[   r\   r]   r^   re   r   r   r_   r   rH   r&   r-   r?   r"   rV   ra   rb   s   @r%   r   r   k   se   ø… ñð$ ƒLØÓØÓØ+/€IÐ(Ó/Ø'+€KÐ$Ó+Ù"'¸Ô"=€J�Ó=Ø€Dˆ-Óôò	
óMògö	r'   N)Ú
__future__r   Údataclassesr   r   Útypingr   r   Útriton.language.corer   Ú+triton.experimental.gluon.language._layoutsr	   Ú__all__r
   r   © r'   r%   ú<module>r}      sl   ðÝ "ç (ß !Ý 5å Ið Øð€ñ �$ÔôX'Ð%ó X'ó ðX'ñv �$ÔôRÐ%ó Ró ñRr'   