ó
    “~iL  ã                  ó    • S SK Jr  S SKJrJr  S SKJrJr  S SKJ	r	  S SK
Jr  SS/r\" SS	9 " S
 S\5      5       r\" SS	9 " S S\5      5       rg)é    )Úannotations)Ú	dataclassÚfield)ÚListÚOptional)Ú_unwrap_if_constexpr)ÚDistributedLayoutÚAMDMFMALayoutÚAMDWMMALayoutT)Úfrozenc                  ó´   ^ • \ rS rSr% SrS\S'   S\S'   S\S'   S\S	'   S
rS\S'   S
rS\S'   \" \	S9r
S\S'   U 4S jrS rSS jrS rS r\S 5       rSrU =r$ )r
   é   aV  
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Úversionú	List[int]Úinstr_shapeÚboolÚ
transposedÚwarps_per_ctaNzOptional[int]Úelement_bitwidthúOptional[List[int]]Útiles_per_warp©Údefault_factoryúList[List[int]]Ú
cga_layoutc                óÊ  >• [         TU ]  S[        U R                  5      5        [         TU ]  S[        U R                  5      5        [         TU ]  S[        U R
                  5      5        [         TU ]  S[        U R                  5      5        [         TU ]  S[        U R                  5      5        [         TU ]  S[        U R                  5      5        U R                  c  [        R                  U SS5        U R                  c.  [        R                  U SS/[        U R                  5      -  5        [        R                  U S	U R                  5        U R                  5         g )
Nr   r   r   r   r   r   é    é   r   )ÚsuperÚ__setattr__r   r   r   r   r   r   r   ÚobjectÚlenr   Úverify)ÚselfÚ	__class__s    €Úl/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/experimental/gluon/language/amd/_layouts.pyÚ__post_init__Ú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à× Ñ Ñ(Ü×Ñ˜tÐ%7¸Ô<Ø×ÑÑ&Ü×Ñ˜tÐ%5¸°s¼SÀ×ASÑASÓ=TÑ7TÔUä×Ñ˜4 ¨t¯©Ô?Ø�‰�ó    c           	     ó¼   • UR                  U R                  U R                  U R                  U R                  U R
                  U R                  U R                  5      $ ©N)Úget_amd_mfma_layoutr   r   r   r   r   r   r   ©r%   Úbuilders     r'   Ú_to_irÚAMDMFMALayout._to_ir<   sN   € Ø×*Ñ*Ø�L‰LØ×ÑØ×ÑØ�O‰OØ�O‰OØ×ÑØ×!Ñ!ó
ð 	
r*   c                ó„  • S nU" U R                   (       a:  U R                    Vs/ s H"  nSR                  [        [        U5      5      PM$     snOS 5      nSU R                   SU" U R
                  5       SU R                   SU" U R                  5       SU R                   SU" U R                  5       SU S3$ s  snf )Nc                óH   • U c  gSR                  [        [        U 5      5      $ ©NÚ Ú_©ÚjoinÚmapÚstr©Úxs    r'   Ú	stringifyÚ'AMDMFMALayout.mangle.<locals>.stringifyI   ó   € Ø‰yØØ—8‘8œC¤ Q›KÓ(Ð(r*   Ú~ÚMFMA_r6   Ú_MFMA)
r   r8   r9   r:   r   r   r   r   r   r   ©r%   r=   Úvecr   s       r'   ÚmangleÚ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   ¥)B=c                ó¬  ^• U R                   S:¼  a  U R                   S::  d   S5       e[        U R                  5      S:X  d   S5       eSS/SS/SS/SS//nU R                  S	S
 U;   d   SU R                   35       eU R                  S;   d   S5       e[        U R                  5      m[        U4S jU R                   5       5      (       d   S5       eg )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   rK   z!element bitwidth must be 32 or 64c              3  ó@   >#   • U  H  n[        U5      T:H  v •  M     g 7fr,   ©r#   ©Ú.0rD   Úranks     €r'   Ú	<genexpr>Ú'AMDMFMALayout.verify.<locals>.<genexpr>Y   ó   øé € Ð?ª¨”3�s“8˜tÖ#ªùó   ƒúcga_layout basis rank mismatch)r   r#   r   r   r   Úallr   )r%   Úvalid_shapesrQ   s     @r'   r$   Ú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  • [        U R                  [        U R                  5      U R                  [        U R
                  5      U R                  (       a  U R                  OS U R                  (       a  [        U R                  5      OS [        S U R                   5       5      45      $ )Nc              3  ó8   #   • U  H  n[        U5      v •  M     g 7fr,   ©Útuple©rP   rD   s     r'   rR   Ú)AMDMFMALayout.__hash__.<locals>.<genexpr>c   ó   é € Ð8ª ”%˜—*�*ªùó   ‚)	Úhashr   r]   r   r   r   r   r   r   ©r%   s    r'   Ú__hash__ÚAMDMFMALayout.__hash__[   s{   € ÜØ�L‰LÜ�$×"Ñ"Ó#Ø�O‰OÜ�$×$Ñ$Ó%Ø%)×%:×%:ˆD×!Ò!ÀØ*.×*=×*=ŒE�$×%Ñ%Ô&À4ÜÑ8¨¯ªÓ8Ó8ð
ó ð 	r*   c                ó,   • [        U R                  5      $ r,   ©r#   r   rc   s    r'   rQ   ÚAMDMFMALayout.rankf   ó   € ä�4×%Ñ%Ó&Ð&r*   © ©Úreturnr:   )Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__Ú__annotations__r   r   r   Úlistr   r(   r0   rE   r$   rd   ÚpropertyrQ   Ú__static_attributes__Ú__classcell__©r&   s   @r'   r
   r
      su   ø‡ ñð& ƒLØÓØÓØÓØ&*Ð�mÓ*Ø*.€NÐ'Ó.Ù"'¸Ñ"=€J�Ó=õò 	
ôAòbò	ð ñ'ó ö'r*   c                  óª   ^ • \ rS rSr% SrS\S'   S\S'   S\S'   S	rS
\S'   S	rS
\S'   \" \	S9r
S\S'   U 4S jrS rSS jrS rS r\S 5       rSrU =r$ )r   ék   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.
    warps_per_cta (List[int]): Number of warps per CTA.
    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.

Current supported versions:

- 1: RDNA3; e.g., gfx1100, gfx1101
- 2: RDNA4; e.g., gfx1200, gfx1201
- 3: gfx1250
r   r   r   r   r   r   Nr   r   r   r   r   r   c                óP  >• [         TU ]  S[        U R                  5      5        [         TU ]  S[        U R                  5      5        [         TU ]  S[        U R
                  5      5        U R                  c  S/[        U R
                  5      -  nO[        U R                  5      n[         TU ]  SU5        U R                  b  [        U R                  5      O/ SQn[         TU ]  S[        U5      5        [        R                  U SU R                  5        U R                  5         g )	Nr   r   r   r   r   )rJ   rJ   rJ   r   r   )r    r!   r   r   r   r   r   r#   r   r"   r   r$   )r%   r   r   r&   s      €r'   r(   ÚAMDWMMALayout.__post_init__„   sæ   ø€ Ü‰Ñ˜IÔ';¸D¿L¹LÓ'IÔJÜ‰Ñ˜LÔ*>¸t¿¹Ó*OÔPÜ‰Ñ˜OÔ-AÀ$×BTÑBTÓ-UÔVà×ÑÑ&Ø˜S¤3 t×'9Ñ'9Ó#:Ñ:‰Nä1°$×2EÑ2EÓFˆNä‰ÑÐ,¨nÔ=à@D×@PÑ@PÑ@\Ô*¨4×+;Ñ+;Ô<ÒbnˆÜ‰Ñ˜MÔ+?ÀÓ+LÔMÜ×Ñ˜4 ¨t¯©Ô?Ø�‰�r*   c                ó¦   • UR                  U R                  U R                  U R                  U R                  U R
                  U R                  5      $ r,   )Úget_amd_wmma_layoutr   r   r   r   r   r   r.   s     r'   r0   ÚAMDWMMALayout._to_ir•   sE   € Ø×*Ñ*Ø�L‰LØ�O‰OØ×ÑØ×ÑØ�O‰OØ×Ñó
ð 	
r*   c                ój  • S nU" U R                   (       a:  U R                    Vs/ s H"  nSR                  [        [        U5      5      PM$     snOS 5      nSU R                   SU R
                   SU" U R                  5       SU" U R                  5       SU" U R                  5       SU S3$ s  snf )Nc                óH   • U c  gSR                  [        [        U 5      5      $ r4   r7   r;   s    r'   r=   Ú'AMDWMMALayout.mangle.<locals>.stringify¡   r?   r*   r@   ÚWMMA_r6   Ú_WMMA)	r   r8   r9   r:   r   r   r   r   r   rC   s       r'   rE   ÚAMDWMMALayout.mangleŸ   sÙ   € ò	)ñ
 ÐW[×Wf×WfÀ4Ç?Â?ÓSÂ?¸C §¡¬¬S°#«Ö 7Á?ÒSÐlpÓqˆ
Ø�t—|‘|�n A d§o¡oÐ%6°a¹	À$×BTÑBTÓ8UÐ7VÐVWÑXaÐbf×buÑbuÓXvÐWwÐwxñ  zCð  DH÷  DTñ  DTó  zUð  yVð  VWð  Xbð  Wcð  chð  ið  	iùò  Ts   ¥)B0c                óÔ   ^• U R                   S:¼  a  U R                   S::  d   S5       e[        U R                  5      m[        U4S jU R                   5       5      (       d   S5       eg )Nr   rI   z#version must be in the [1, 3] rangec              3  ó@   >#   • U  H  n[        U5      T:H  v •  M     g 7fr,   rN   rO   s     €r'   rR   Ú'AMDWMMALayout.verify.<locals>.<genexpr>­   rT   rU   rV   )r   r#   r   rW   r   )r%   rQ   s    @r'   r$   ÚAMDWMMALayout.verify©   sW   ø€ Ø�|‰|˜qÓ  T§\¡\°QÓ%6Ð]Ð8]Ó]Ð6ä�4×%Ñ%Ó&ˆÜÔ?¨t¯ªÓ?×?Ñ?ÐaÐAaÓaÑ?r*   c                ó>  • [        U R                  U R                  [        U R                  5      U R
                  (       a  [        U R
                  5      OS U R                  (       a  [        U R                  5      OS [        S U R                   5       5      45      $ )Nc              3  ó8   #   • U  H  n[        U5      v •  M     g 7fr,   r\   r^   s     r'   rR   Ú)AMDWMMALayout.__hash__.<locals>.<genexpr>¶   r`   ra   )rb   r   r   r]   r   r   r   r   rc   s    r'   rd   ÚAMDWMMALayout.__hash__¯   sr   € ÜØ�L‰LØ�O‰OÜ�$×$Ñ$Ó%Ø*.×*=×*=ŒE�$×%Ñ%Ô&À4Ø'+×'7×'7ŒE�$×"Ñ"Ô#¸TÜÑ8¨¯ªÓ8Ó8ð
ó ð 	r*   c                ó,   • [        U R                  5      $ r,   rg   rc   s    r'   rQ   ÚAMDWMMALayout.rank¹   ri   r*   rj   rk   )rm   rn   ro   rp   rq   rr   r   r   r   rs   r   r(   r0   rE   r$   rd   rt   rQ   ru   rv   rw   s   @r'   r   r   k   so   ø‡ ñð  ƒLØÓØÓØ'+€KÐ$Ó+Ø*.€NÐ'Ó.Ù"'¸Ñ"=€J�Ó=õò"
ôiòbòð ñ'ó ö'r*   N)Ú
__future__r   Údataclassesr   r   Útypingr   r   Útriton.language.corer   Ú+triton.experimental.gluon.language._layoutsr	   Ú__all__r
   r   rj   r*   r'   Ú<module>r•      sl   ðÝ "ç (ß !Ý 5å Ið Øð€ñ �$ÑôX'Ð%ó X'ó ðX'ñv �$ÑôO'Ð%ó O'ó ñO'r*   