o
    h#                     @   s   d dl mZ d dlmZmZ d dlmZmZ g dZdd Z	G dd dZ
ed	d
G dd de
Zed	d
G dd de
Zed	d
G dd de
ZG dd dZed	d
G dd deZed	d	dG dd deZdS )    )	dataclass)ListOptional)_unwrap_if_constexpr_unwrap_shape)BlockedLayoutSliceLayoutDistributedLinearLayoutNVMMASharedLayoutSwizzledSharedLayoutc                 C   s:   |pdg|  }|pdg|  }|pt tt| }|||fS )N   )listreversedrange)rankctas_per_cgacta_split_num	cta_order r   f/var/www/html/scripts/venv/lib/python3.10/site-packages/triton/experimental/gluon/language/_layouts.py_realize_cta_layout   s   
r   c                   @      e Zd ZdS )DistributedLayoutN__name__
__module____qualname__r   r   r   r   r          r   T)frozenc                       s   e Zd ZU ee ed< ee ed< ee ed< ee ed< dZeee  ed< dZeee  ed< dZ	eee  ed<  fd	d
Z
dd ZdefddZ  ZS )r   size_per_threadthreads_per_warpwarps_per_ctaorderNr   r   r   c                    s(  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	 t
| j}t
| j|ksTJ t
| j|ks]J t
| j|ksfJ | jd u stt
| j|kstJ | jd u st
| j|ksJ | j	d u st
| j	|ksJ d S d S )Nr   r    r!   r"   r   r   r   )super__setattr__r   r   r    r!   r"   r   r   r   lenselfr   	__class__r   r   __post_init__#   s   
$zBlockedLayout.__post_init__c              	   C   sB   t | j}t|| j| j| j\}}}|| j| j| j| j	|||S N)
r%   r   r   r   r   r   get_blocked_layoutr    r!   r"   r'   builderr   r   r   r   r   r   r   _to_ir4   s   

zBlockedLayout._to_irreturnc           	      C   s~   dd }|| j }|| j}|| j}|| j}|| j}|| j}|| j}d| d| d| d| d| d| d| dS )Nc                 S      | d u rdS d tt| S N _joinmapstrxr   r   r   	stringifyD      z'BlockedLayout.mangle.<locals>.stringifyB)r   r    r!   r"   r   r   r   )	r'   r;   r   r    r!   r"   r   r   r   r   r   r   mangleB   s   






0zBlockedLayout.mangle)r   r   r   r   int__annotations__r   r   r   r   r*   r/   r8   r>   __classcell__r   r   r(   r   r      s   
 r   c                       sD   e Zd ZU eed< eed<  fddZdd Zdefdd	Z	  Z
S )
r   dimparentc                    s,   t  dt| j t  dt| j d S )NrB   rC   )r#   r$   r   rB   rC   r'   r(   r   r   r*   X   s   zSliceLayout.__post_init__c                 C   s   | | j| j|S r+   )get_slice_layoutrB   rC   r/   r'   r.   r   r   r   r/   \   s   
zSliceLayout._to_irr0   c                 C   s   d| j  d| j  dS )NSLr4   )rB   rC   r>   rD   r   r   r   r>   b   s   zSliceLayout.mangle)r   r   r   r?   r@   r   r*   r/   r8   r>   rA   r   r   r(   r   r   S   s   
 r   c                       sz   e Zd ZU 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dZdd	 Zd
d Z  Z	S )r	   	reg_bases
lane_bases
warp_basesblock_basesshapec                    s   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sDJ q:| jD ]
}t||ksRJ qH| jD ]
}t||ks`J qV| jD ]
}t||ksnJ qdd S )NrH   rI   rJ   rK   rL   )	r#   r$   r   rH   rI   rJ   rK   rL   r%   )r'   r   basisr(   r   r   r*   n   s   




z%DistributedLinearLayout.__post_init__c                 C   s   | | j| j| j| j| jS r+   )get_distributed_linear_layoutrH   rI   rJ   rK   rL   rF   r   r   r   r/      s   zDistributedLinearLayout._to_irc                 C   s.   d| j  d| j d| j d| j d| j dS )NDLLr4   )rH   rI   rJ   rK   rL   rD   r   r   r   r>      s   .zDistributedLinearLayout.mangle)
r   r   r   r   r?   r@   r*   r/   r>   rA   r   r   r(   r   r	   f   s   
 r	   c                   @   r   )SharedLayoutNr   r   r   r   r   rP      r   rP   c                       s   e Zd ZU eed< eed< eed< dZeed< dZeed< dZe	e
e  ed< dZe	e
e  ed	< dZe	e
e  ed
<  fddZdd ZdefddZ  ZS )r
   swizzle_byte_widthelement_bitwidthr   F
transposed
fp4_paddedNr   r   r   c                    s  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	 t  dt| j
 | jd	v sWJ | jd
v s^J | j}| jd u sot| j|ksoJ | j	d u s}t| j	|ks}J | j
d u st| j
|ksJ d S d S )NrQ   rR   r   rS   rT   r   r   r   )          @   )r   rW   rX      )r#   r$   r   rQ   rR   r   rS   rT   r   r   r   r%   r&   r(   r   r   r*      s   $zNVMMASharedLayout.__post_init__c              	   C   s:   t | j| j| j| j\}}}|| j| j| j| j	|||S r+   )
r   r   r   r   r   get_nvmma_shared_layoutrQ   rR   rS   rT   )r'   r.   r   r   r   r   r   r   r/      s   
zNVMMASharedLayout._to_irr0   c              	   C   s&   d| j  d| j d| j d| j d	S )NNVMMA_r4   _NVMMA)rQ   rR   rS   rT   rD   r   r   r   r>      s   &zNVMMASharedLayout.mangle)r   r   r   r?   r@   rS   boolrT   r   r   r   r   r   r*   r/   r8   r>   rA   r   r   r(   r   r
      s   
 r
   )r   eqc                       s   e Zd ZU eed< eed< eed< ee ed< dZeee  ed< dZeee  ed< dZ	eee  ed<  fd	d
Z
dd ZdefddZ  ZS )r   vec	per_phase	max_phaser"   Nr   r   r   c                    s   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	 t
| j}| jd u sYt
| j|ksYJ | jd u sgt
| j|ksgJ | j	d u sut
| j	|kswJ d S d S )Nr_   r`   ra   r"   r   r   r   )r#   r$   r   r_   r`   ra   r"   r   r   r   r%   r&   r(   r   r   r*      s   
$z"SwizzledSharedLayout.__post_init__c              	   C   sN   t | j}t|| j| j| j\}}}|t| jt| j	t| j
| j|||S r+   )r%   r"   r   r   r   r   get_swizzled_shared_layoutr   r_   r`   ra   r-   r   r   r   r/      s   

zSwizzledSharedLayout._to_irr0   c                 C   sV   dd }d| j  d| j d| j d|| j d|| j d|| j d|| j dS )Nc                 S   r1   r2   r5   r9   r   r   r   r;      r<   z.SwizzledSharedLayout.mangle.<locals>.stringifySSS_r4   _SSS)r_   r`   ra   r"   r   r   r   )r'   r;   r   r   r   r>      s   NzSwizzledSharedLayout.mangle)r   r   r   r?   r@   r   r   r   r   r   r*   r/   r8   r>   rA   r   r   r(   r   r      s   
 r   N)dataclassesr   typingr   r   triton.language.corer   r   __all__r   r   r   r   r	   rP   r
   r   r   r   r   r   <module>   s"    	9!
,