a
    kh#                     @   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   Y/var/www/auris/lib/python3.9/site-packages/triton/experimental/gluon/language/_layouts.py_realize_cta_layout   s    r   c                   @   s   e Zd ZdS )DistributedLayoutN__name__
__module____qualname__r   r   r   r   r      s   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ed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sJ t
| j|ksJ t
| j|ksJ | jd u st
| j|ksJ | jd u st
| j|ksJ | j	d u s(t
| j	|ks(J 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   Z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   s   | d u rdS d tt| S N _joinmapstrxr   r   r   	stringifyD   s    z'BlockedLayout.mangle.<locals>.stringifyB)r   r   r   r    r   r   r   )	r%   r8   r   r   r   r    r   r   r   r   r   r   mangleB   s    






zBlockedLayout.mangle)r   r   r   r   int__annotations__r   r   r   r   r(   r,   r5   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eddd	Z	  Z
S )
r   dimparentc                    s,   t  dt| j t  dt| j d S )Nr>   r?   )r!   r"   r   r>   r?   r%   r&   r   r   r(   X   s    zSliceLayout.__post_init__c                 C   s   | | j| j|S r)   )Zget_slice_layoutr>   r?   r,   r%   r+   r   r   r   r,   \   s    
zSliceLayout._to_irr-   c                 C   s   d| j  d| j  dS )NZSLr1   )r>   r?   r:   r@   r   r   r   r:   b   s    zSliceLayout.mangle)r   r   r   r;   r<   r   r(   r,   r5   r:   r=   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stJ qt| jD ]}t||ksJ q| jD ]}t||ksJ q| jD ]}t||ksJ qd S )NrB   rC   rD   rE   rF   )	r!   r"   r   rB   rC   rD   rE   rF   r#   )r%   r   Zbasisr&   r   r   r(   n   s    




z%DistributedLinearLayout.__post_init__c                 C   s   | | j| j| j| j| jS r)   )Zget_distributed_linear_layoutrB   rC   rD   rE   rF   rA   r   r   r   r,      s    zDistributedLinearLayout._to_irc                 C   s.   d| j  d| j d| j d| j d| j dS )NZDLLr1   )rB   rC   rD   rE   rF   r@   r   r   r   r:      s    zDistributedLinearLayout.mangle)
r   r   r   r   r;   r<   r(   r,   r:   r=   r   r   r&   r   r	   f   s   
r	   c                   @   s   e Zd ZdS )SharedLayoutNr   r   r   r   r   rG      s   rG   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ed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 sJ | jd
v sJ | j}| jd u st| j|ksJ | j	d u st| j	|ksJ | j
d u st| j
|ksJ d S )NrH   rI   r   rJ   rK   r   r   r   )          @   )r   rN   rO      )r!   r"   r   rH   rI   r   rJ   rK   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   Zget_nvmma_shared_layoutrH   rI   rJ   rK   )r%   r+   r   r   r   r   r   r   r,      s    
zNVMMASharedLayout._to_irr-   c              	   C   s&   d| j  d| j d| j d| j d	S )NZNVMMA_r1   Z_NVMMA)rH   rI   rJ   rK   r@   r   r   r   r:      s    zNVMMASharedLayout.mangle)r   r   r   r;   r<   rJ   boolrK   r   r   r   r   r   r(   r,   r5   r:   r=   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ed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 st
| j|ksJ | jd u st
| j|ksJ | j	d u st
| j	|ksJ d S )NrS   rT   rU   r    r   r   r   )r!   r"   r   rS   rT   rU   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   Zget_swizzled_shared_layoutr   rS   rT   rU   r*   r   r   r   r,      s    

zSwizzledSharedLayout._to_irr-   c                 C   sV   dd }d| j  d| j d| j d|| j d|| j d|| j d|| j dS )Nc                 S   s   | d u rdS d tt| S r/   r2   r6   r   r   r   r8      s    z.SwizzledSharedLayout.mangle.<locals>.stringifyZSSS_r1   Z_SSS)rS   rT   rU   r    r   r   r   )r%   r8   r   r   r   r:      s    zSwizzledSharedLayout.mangle)r   r   r   r;   r<   r   r   r   r   r   r(   r,   r5   r:   r=   r   r   r&   r   r      s   
r   N)Zdataclassesr   typingr   r   Ztriton.language.corer   r   __all__r   r   r   r   r	   rG   r
   r   r   r   r   r   <module>   s    	9!,
