o
    i?                     @  s   d dl mZ d dlmZmZmZ d dlmZ d dlm	Z	m
Z
 d dlm  m  m  mZ d dlmZ d dlmZmZ erCd dlmZ g d	Zed
dG dd de	ZG dd de
Zed#ddZed$ddZed$ddZe		d%d&d!d"ZdS )'    )annotations)ListTupleTYPE_CHECKING)	dataclass)	base_type
base_valueN)NVMMASharedLayout)builtin_unwrap_if_constexpr)ir)async_copy_global_to_sharedasync_copy_shared_to_global
store_waitT)eqc                   @  s`   e Zd ZU ded< ded< ded< ded< d d
dZd!ddZd"ddZd#ddZd ddZdS )$tensor_descriptor_typettgl.block_type
block_typezttgl.tuple_type
shape_typestrides_typer	   layoutreturnstrc                 C  s   d| j  d| j dS )Nztensor_descriptor<z, >)r   r   self r   g/home/kuhnn/.local/lib/python3.10/site-packages/triton/experimental/gluon/language/nvidia/hopper/tma.py__str__   s   ztensor_descriptor_type.__str__builder
ir.builderir.typec                 C  s*   | j j }|| j ||| j|S N)r   
element_tyis_int_signed!get_tensor_descriptor_layout_typeto_irr   _to_ir)r   r   	is_signedr   r   r   r'      s   

ztensor_descriptor_type._to_irhandlesList[ir.value]cursorintTuple[tensor_descriptor, int]c                 C  sR   || }|d7 }| j ||\}}| j||\}}t|||| j| jd}||fS )N   )r   )r   _unflatten_irr   tensor_descriptorr   r   )r   r)   r+   handleshapestridesvaluer   r   r   r/   !   s   z$tensor_descriptor_type._unflatten_iroutList[ir.type]Nonec                 C  sT   | j j }|| j ||| j|}|| | j	|| | j
	|| d S r"   )r   r#   r$   r%   r&   r   r'   appendr   _flatten_ir_typesr   )r   r   r5   r(   tyr   r   r   r9   )   s   


z(tensor_descriptor_type._flatten_ir_typesc                 C  s   d| j   d| j  dS )NTD_)r   mangler   r   r   r   r   r=   4   s   ztensor_descriptor_type.mangleN)r   r   )r   r    r   r!   )r)   r*   r+   r,   r   r-   )r   r    r5   r6   r   r7   )	__name__
__module____qualname____annotations__r   r'   r/   r9   r=   r   r   r   r   r      s   
 



r   c                   @  sP   e Zd Zddd	ZdddZedd Zedd Zedd Zedd Z	dS )r0   r2   List[ttgl.tensor]r3   r   r   r   r	   c                 C  s<   || _ t|| _t|| _t|| jj| jj|d| _d S )N)r   r   r   )r1   ttgltupler2   r3   r   type)r   r1   r2   r3   r   r   r   r   r   __init__:   s   ztensor_descriptor.__init__r)   r*   r   r7   c                 C  s(   | | j | j| | j| d S r"   )r8   r1   r2   _flatten_irr3   )r   r)   r   r   r   rG   B   s   ztensor_descriptor._flatten_irc                 C     | j jS r"   )rE   r   r   r   r   r   r   G      ztensor_descriptor.block_typec                 C  
   | j jjS r"   )rE   r   r2   r   r   r   r   block_shapeK      
ztensor_descriptor.block_shapec                 C  rJ   r"   )rE   r   r#   r   r   r   r   dtypeO   rL   ztensor_descriptor.dtypec                 C  rH   r"   )rE   r   r   r   r   r   r   S   rI   ztensor_descriptor.layoutN)r2   rB   r3   rB   r   r   r   r	   )r)   r*   r   r7   )
r>   r?   r@   rF   rG   propertyr   rK   rM   r   r   r   r   r   r0   8   s    




r0   c                 C  s8   |j |dd}||}|j| j||j|j|j d S NF)require_i64)_convert_to_ir_values	to_tensorr   %create_async_tma_copy_global_to_localr1   )tensor_desccoordbarrierresultpred	_semanticr   r   r   r   X   s
   
r   c                 C  s&   |j |dd}|j| j||j d S rO   )rQ   r   %create_async_tma_copy_local_to_globalr1   )rT   rU   srcrY   r   r   r   r   `   s   r   c                 C  s   t | } |j|  d S r"   )r   r   create_async_tma_store_wait)pendingsrY   r   r   r   r   f   s   r   zerobasettgl.tensorr2   rB   r3   rK   List[ttgl.constexpr]r   r	   r   c                   s  t |}t |}t|}d|  krdksn td| dt||kr1td| dt| t||krCtd| dt| t| jtjsLJ | jjjd	 }t |d
 }	|	| dk rptd|	 d| d|	|  dt |d
 }
|
dkrtd|
  fdd|D } fdd|D }t	|}t| j
tjsJ t| j
j|}| j} |}t |}t|tsJ dt|j
}t|j
}t||||}| j
j r|tjjjkrtd j| j|dd |D dd |D |}t|||||S )Nr.      z Expected 1 <= ndim <= 5 but got z dimensionsz	Expected z strides but got zExpected block_shape to have z dimensions but got       zRDescriptor block shape must have at least 16 bytes in the last dimension, but got z * z = z bytesz-Tensor descriptor last dim must be 1 but got c                   s   g | ]	}  |tjqS r   )make_scalarrC   int32.0xrY   r   r   
<listcomp>   s    z*make_tensor_descriptor.<locals>.<listcomp>c                   s    g | ]}  t|tjqS r   )rf   rC   r   int64rh   rk   r   r   rl      s     z)Expected layout to be a NVMMASharedLayoutz8Padding option `nan` is not supported for integer blocksc                 S     g | ]}|j qS r   r1   ri   sr   r   r   rl          c                 S  rn   r   ro   rp   r   r   r   rl      rr   )r   len
ValueError
isinstancerM   rC   pointer_typer#   primitive_bitwidth_unwrap_shaperE   r   r1   _str_to_padding_optionr	   rD   r   is_intr   PADDING_OPTIONPAD_NANr   create_make_tensor_descriptorr'   r0   )r_   r2   r3   rK   r   padding_optionrY   ndim	elem_sizecontig_dim_sizelast_strider   base_handlepaddingr   r   r:   r1   r   rk   r   make_tensor_descriptorl   sV   



r   )TNr"   )r^   N)r_   r`   r2   rB   r3   rB   rK   ra   r   r	   r   r0   )
__future__r   typingr   r   r   dataclassesr   triton.language.corer   r   (triton.experimental.gluon.language._coreexperimentalgluonlanguage_corerC   +triton.experimental.gluon.language._layoutsr	   r
   r   	triton._Cr   __all__r   r0   r   r   r   r   r   r   r   r   <module>   s.    ( 