
    rh                        d dl mZ d dlmZmZmZ d dlmZ d dlm	c m
c mc mZ d dlmZ d dlmZmZ erd dlmZ g dZ ed	
       G d d             Z G d d      Zedd       Zedd       Zedd       Zy)    )annotations)ListTupleTYPE_CHECKING)	dataclassN)NVMMASharedLayout)builtin_unwrap_if_constexpr)ir)async_copy_global_to_sharedasync_copy_shared_to_global
store_waitT)eqc                  V    e Zd ZU ded<   ded<   ded<   ded<   ddZdd	Zdd
ZddZy)tensor_descriptor_typettgl.block_type
block_typezttgl.tuple_type
shape_typestrides_typer   layoutc                <    d| j                    d| j                   dS )Nztensor_descriptor<z, >)r   r   selfs    /var/www/html/ai-insurance-compliance-backend/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/nvidia/hopper/tma.py__str__ztensor_descriptor_type.__str__   s     #DOO#4Bt{{m1EE    c                    ||   }|dz  }| j                   j                  ||      \  }}| j                  j                  ||      \  }}t        |||| j                  | j
                        }||fS )N   )r   )r   _unflatten_irr   tensor_descriptorr   r   )r   handlescursorhandleshapestridesvalues          r   r    z$tensor_descriptor_type._unflatten_ir   sq    !55gvFv++99'6J!&%$//RVR]R]^f}r   c                h   | j                   j                  j                         }|j                  | j                   j	                  |      || j
                  j                  |            }|j                  |       | j                  j                  ||       | j                  j                  ||       y N)r   
element_tyis_int_signed!get_tensor_descriptor_layout_typeto_irr   _to_irappendr   _flatten_ir_typesr   )r   builderout	is_signedtys        r   r0   z(tensor_descriptor_type._flatten_ir_types    s    OO..<<>	66OO!!'*KKw'

 	

2))'37++GS9r   c                l    d| j                   j                   d| j                  j                          dS )NTD_)r   mangler   r   s    r   r8   ztensor_descriptor_type.mangle+   s0    DOO**+1T[[-?-?-A,B"EEr   N)returnstr)r"   List[ir.value]r#   intr9   zTuple[tensor_descriptor, int])r1   z
ir.builderr2   zList[ir.type]r9   None)__name__
__module____qualname____annotations__r   r    r0   r8    r   r   r   r      s0    !!F	:Fr   r   c                  `    e Zd Z	 	 ddZd	dZed        Zed        Zed        Zed        Z	y)
r!   c                    || _         t        j                  |      | _        t        j                  |      | _        t        || j                  j                  | j                  j                  |      | _        y )N)r   r   r   )r$   ttgltupler%   r&   r   type)r   r$   r%   r&   r   r   s         r   __init__ztensor_descriptor.__init__1   sS    ZZ&
zz'**:$**//`d`l`l`q`q28:	r   c                    |j                  | j                         | j                  j                  |       | j                  j                  |       y r)   )r/   r$   r%   _flatten_irr&   )r   r"   s     r   rJ   ztensor_descriptor._flatten_ir9   s6    t{{#

w'  )r   c                .    | j                   j                  S r)   )rG   r   r   s    r   r   ztensor_descriptor.block_type>   s    yy###r   c                B    | j                   j                  j                  S r)   )rG   r   r%   r   s    r   block_shapeztensor_descriptor.block_shapeB   s    yy##)))r   c                B    | j                   j                  j                  S r)   )rG   r   r*   r   s    r   dtypeztensor_descriptor.dtypeF   s    yy##...r   c                .    | j                   j                  S r)   )rG   r   r   s    r   r   ztensor_descriptor.layoutJ   s    yyr   N)r%   List[ttgl.tensor]r&   rQ   r   r   r   r   )r"   r;   r9   r=   )
r>   r?   r@   rH   rJ   propertyr   rM   rO   r   rB   r   r   r!   r!   /   sd    :*:*
 $ $ * * / /    r   r!   c                    |j                  |d      }|j                  |      }|j                  j                  | j                  ||j                  |j                  |j                         y NF)require_i64)_convert_to_ir_values	to_tensorr1   %create_async_tma_copy_global_to_localr$   )tensor_desccoordbarrierresultpred	_semantics         r   r   r   O   s`    ++Eu+EEt$D;;K<N<NPUW^WeWegmgtgt<@KKIr   c                    |j                  |d      }|j                  j                  | j                  ||j                         y rT   )rV   r1   %create_async_tma_copy_local_to_globalr$   )rY   rZ   srcr^   s       r   r   r   W   s=    ++Eu+EE;;K<N<NPUWZWaWabr   c                P    t        |       } |j                  j                  |        y r)   )r
   r1   create_async_tma_store_wait)pendingsr^   s     r   r   r   ]   s     #H-H11(;r   )TNr)   )
__future__r   typingr   r   r   dataclassesr   (triton.experimental.gluon.language._coreexperimentalgluonlanguage_corerE   +triton.experimental.gluon.language._layoutsr   r	   r
   	triton._Cr   __all__r   r!   r   r   r   rB   r   r   <module>rp      s    " - - ! 7 7 I R
V dF F F@   @ 	I 	I 	c 	c
 	< 	<r   