
    xpi                       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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e	             Z G d de
      Zedd       Zedd       Zedd       Zy)    )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                  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    v/opt/services/ai/voice_agent/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   r2   z(tensor_descriptor_type._flatten_ir_types!   s    OO..<<>	66OO!!'*KKw'

 	

2))'37++GS9r   c                t    d| j                   j                          d| j                  j                          dS )NTD_)r   mangler   r   s    r   r:   ztensor_descriptor_type.mangle,   s3    DOO**,-Qt{{/A/A/C.DBGGr   N)returnstr)r$   List[ir.value]r%   intr;   zTuple[tensor_descriptor, int])r3   z
ir.builderr4   zList[ir.type]r;   None)__name__
__module____qualname____annotations__r   r"   r2   r:    r   r   r   r      s0    !!F	:H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__2   sS    ZZ&
zz'**:$**//`d`l`l`q`q28:	r   c                    |j                  | j                         | j                  j                  |       | j                  j                  |       y r+   )r1   r&   r'   _flatten_irr(   )r   r$   s     r   rL   ztensor_descriptor._flatten_ir:   s6    t{{#

w'  )r   c                .    | j                   j                  S r+   )rI   r   r   s    r   r   ztensor_descriptor.block_type?   s    yy###r   c                B    | j                   j                  j                  S r+   )rI   r   r'   r   s    r   block_shapeztensor_descriptor.block_shapeC   s    yy##)))r   c                B    | j                   j                  j                  S r+   )rI   r   r,   r   s    r   dtypeztensor_descriptor.dtypeG   s    yy##...r   c                .    | j                   j                  S r+   )rI   r   r   s    r   r   ztensor_descriptor.layoutK   s    yyr   N)r'   List[ttgl.tensor]r(   rS   r   r   r   r
   )r$   r=   r;   r?   )
r@   rA   rB   rJ   rL   propertyr   rO   rQ   r   rD   r   r   r#   r#   0   sd    :*:*
 $ $ * * / /    r   r#   c                    |j                  |d      }|j                  |      }|j                  j                  | j                  ||j                  |j                  |j                         y NF)require_i64)_convert_to_ir_values	to_tensorr3   %create_async_tma_copy_global_to_localr&   )tensor_desccoordbarrierresultpred	_semantics         r   r   r   P   s`    ++Eu+EEt$D;;K<N<NPUW^WeWegmgtgt<@KKIr   c                    |j                  |d      }|j                  j                  | j                  ||j                         y rV   )rX   r3   %create_async_tma_copy_local_to_globalr&   )r[   r\   srcr`   s       r   r   r   X   s=    ++Eu+EE;;K<N<NPUWZWaWabr   c                P    t        |       } |j                  j                  |        y r+   )r   r3   create_async_tma_store_wait)pendingsr`   s     r   r   r   ^   s     #H-H11(;r   )TNr+   )
__future__r   typingr   r   r   dataclassesr   triton.language.corer   r	   (triton.experimental.gluon.language._coreexperimentalgluonlanguage_corerG   +triton.experimental.gluon.language._layoutsr
   r   r   	triton._Cr   __all__r   r#   r   r   r   rD   r   r   <module>rs      s    " - - ! 6 7 7 I R
V dHY H H@ 
  @ 	I 	I 	c 	c
 	< 	<r   