
    ^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';DLL'IJM+?@P@P+QRL*>t*OPO-A$BTBT-UV.0DTEZEZ0[\,.B4CVCV.WXL*>t*OP  (G 2B7&G 01#D<N<N8O2O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    **LLOOOO!!
 	
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88CQK((r'   ~MFMA_r2   _MFMA)
r   r4   r5   r6   r   r   r   r   r   r   )r#   r9   vecr   s       r%   manglezAMDMFMALayout.mangleG   s    	)
 W[WfWf4??SCS# 7Slpq
t||nAi0@0@&A%B!DOOCTTUV_`d`r`rVsUttuvz  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     ?3s8t#?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2r(RGaW=!$4c8PQUQaQaPb6cc4$$0U2UU04%%&?t??aAa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{    LL$""#OO$$$%%)%:%: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';DLL'IJL*>t*OPLDXY]YhYhDi*j54;*jkKSWSaSaSm*>t~~*NP!& "%[ Psu	w AE@P@P@\*4+;+;<bnM+?+LML*>t*OP26))2G#DII.QFD) +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    **LLOOOONNOOII
 	
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#chhs3}5CSSd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 &doo6
$T^^4	%doo6
t||nAdoo%6a
|1YKqQZ[_[k[kQlPmmnoynzz{  }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     DS3s8tyy(Ds   "%rJ   )r   r!   rd   rH   rK   r   rU   s   `r%   r"   zAMDWMMALayout.verify   s{    ||q T\\Q%6]8]]6t!#tq)*dii7Y9YY7DDOODDfFf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%*7rS   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    LLOO888777'+'7'7E$""#T888II
  	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'   