
    WiL                        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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         |   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        j                  | dd       | j                  .t        j                  | ddgt        | j                        z         t        j                  | d	| j                         | j                          y )
Nr   r   r   r   r   r          r   )super__setattr__r   r   r   r   r   r   r   objectlenr   verify)self	__class__s    {/home/sietch6/trending-topics-pipeline/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  (t%7<&t%5sSASAS=T7TU4t?    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_r3   _MFMA)
r   r5   r6   r7   r   r   r   r   r   r   r$   r:   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   rF   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        ?3s8t#?   cga_layout basis rank mismatch)r   r"   r   r   r   allr   )r$   valid_shapesrL   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rK   r@   s     r&   rM   z)AMDMFMALayout.__hash__.<locals>.<genexpr>c        8%*8   )	hashr   rV   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   r[   s    r&   rL   zAMDMFMALayout.rankf       4%%&&r(   returnr7   )__name__
__module____qualname____doc____annotations__r   r   r   listr   r'   r.   rA   r#   r\   propertyrL   __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<    fdZ
d ZddZd Zd Zed        Z xZS )r   a8  
    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                L   t         |   dt        | j                               t         |   dt        | j                               t         |   dt        | j
                               | j                  dgt        | j
                        z  }nt        | j                        }t         |   d|       | j                  t        | j                        ng d}t         |   dt        |             t        j                  | d| j                         | j                          y )	Nr   r   r   r   r   )rE   rE   rE   r   r   )r   r    r   r   r   r   r   r"   r   r!   r   r#   )r$   r   r   r%   s      r&   r'   zAMDWMMALayout.__post_init__   s    I';DLL'IJL*>t*OPO-A$BTBT-UV&S3t'9'9#::N1$2E2EFN,n=@D@P@P@\*4+;+;<bnM+?+LM4t?r(   c                    |j                  | j                  | j                  | j                  | j                  | j
                  | j                        S r*   )get_amd_wmma_layoutr   r   r   r   r   r   r,   s     r&   r.   zAMDWMMALayout._to_ir   sE    **LLOOOO
 	
r(   c                Z   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| dS c c}w )Nc                F    | ydj                  t        t        |             S r1   r4   r8   s    r&   r:   z'AMDWMMALayout.mangle.<locals>.stringify   r;   r(   r<   WMMA_r3   _WMMA)	r   r5   r6   r7   r   r   r   r   r   r?   s       r&   rA   zAMDWMMALayout.mangle   s    	)
 W[WfWf4??SCS# 7Slpq
t||nAdoo%6a	$BTBT8U7VVWXabfbubuXvWwwx  zC  DH  DT  DT  zU  yV  VW  Xb  Wc  ch  i  	i  Ts   &B(c                    | j                   dk\  r| j                   dk  sJ d       t        | j                        t        fd| j                  D              sJ d       y )Nr   rD   z#version must be in the [1, 3] rangec              3  :   K   | ]  }t        |      k(    y wr*   rI   rJ   s     r&   rM   z'AMDWMMALayout.verify.<locals>.<genexpr>   rN   rO   rP   )r   r"   r   rQ   r   )r$   rL   s    @r&   r#   zAMDWMMALayout.verify   sT    ||q T\\Q%6]8]]64%%&?t??aAaa?r(   c                *   t        | j                  | j                  t        | j                        | j
                  rt        | j
                        nd | j                  rt        | j                        nd t        d | j                  D              f      S )Nc              3  2   K   | ]  }t        |        y wr*   rU   rW   s     r&   rM   z)AMDWMMALayout.__hash__.<locals>.<genexpr>   rX   rY   )rZ   r   r   rV   r   r   r   r   r[   s    r&   r\   zAMDWMMALayout.__hash__   sr    LLOO$$$%*.*=*=E$%%&4'+'7'7E$""#T888
  	r(   c                ,    t        | j                        S r*   r^   r[   s    r&   rL   zAMDWMMALayout.rank   r_   r(   r`   )rb   rc   rd   re   rf   r   r   r   rg   r   r'   r.   rA   r#   r\   rh   rL   ri   rj   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    r(   r&   <module>r      sl    " ( ! 5 I  $X'% X' X'v $O'% O' O'r(   