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   q/home/www/facesmatcher.com/frenv_anti/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   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      | 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.   r7   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 )NrA   rB   )r#   r$   r   rA   rB   r'   r(   r   r   r*   X   s   zSliceLayout.__post_init__c                 C   s   | | j| j|S r+   )Zget_slice_layoutrA   rB   r.   r'   r-   r   r   r   r.   \   s   
zSliceLayout._to_irr/   c                 C   s   d| j  d| j  dS )NZSLr3   )rA   rB   r=   rC   r   r   r   r=   b   s   zSliceLayout.mangle)r   r   r   r>   r?   r   r*   r.   r7   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sDJ q:| jD ]
}t||ksRJ qH| jD ]
}t||ks`J qV| jD ]
}t||ksnJ qdd S )NrE   rF   rG   rH   rI   )	r#   r$   r   rE   rF   rG   rH   rI   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_layoutrE   rF   rG   rH   rI   rD   r   r   r   r.      s   zDistributedLinearLayout._to_irc                 C   s.   d| j  d| j d| j d| j d| j dS )NZDLLr3   )rE   rF   rG   rH   rI   rC   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                   @   r   )SharedLayoutNr   r   r   r   r   rJ      r   rJ   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 )NrK   rL   r   rM   rN   r   r   r   )          @   )r   rQ   rR      )r#   r$   r   rK   rL   r   rM   rN   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_layoutrK   rL   rM   rN   )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_r3   Z_NVMMA)rK   rL   rM   rN   rC   r   r   r   r=      s   &zNVMMASharedLayout.mangle)r   r   r   r>   r?   rM   boolrN   r   r   r   r   r   r*   r.   r7   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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 )NrV   rW   rX   r"   r   r   r   )r#   r$   r   rV   rW   rX   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   rV   rW   rX   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   r0   r1   r4   r8   r   r   r   r:      r;   z.SwizzledSharedLayout.mangle.<locals>.stringifyZSSS_r3   Z_SSS)rV   rW   rX   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.   r7   r=   r@   r   r   r(   r   r      s   
 r   N)dataclassesr   typingr   r   Ztriton.language.corer   r   __all__r   r   r   r   r	   rJ   r
   r   r   r   r   r   <module>   s"    	9!
,