
    h                       S SK Jr  S SKJrJrJr  S SKJr  S SKJ	s  J
s  Js  Jr  S SKJr  S SKJrJr  \(       a  S SKJr  / SQr\" S	S
9 " S S5      5       r " S S5      r\SS j5       r\SS j5       r\SS j5       rg)    )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                  f    \ rS rSr% S\S'   S\S'   S\S'   S\S'   SS	 jrSS
 jrSS jrSS jrSr	g)tensor_descriptor_type   ttgl.block_type
block_typezttgl.tuple_type
shape_typestrides_typer   layoutc                <    SU R                    SU R                   S3$ )Nztensor_descriptor<z, >)r   r   selfs    n/var/www/fran/franai/venv/lib/python3.13/site-packages/triton/experimental/gluon/language/nvidia/hopper/tma.py__str__tensor_descriptor_type.__str__   s     #DOO#4Bt{{m1EE    c                    X   nUS-  nU R                   R                  X5      u  pBU R                  R                  X5      u  pR[        X4XPR                  U R
                  S9nXb4$ )N   )r   )r   _unflatten_irr   tensor_descriptorr   r   )r   handlescursorhandleshapestridesvalues          r   r"   $tensor_descriptor_type._unflatten_ir   sa    !55gF++99'J!&//RVR]R]^}r   c                d   U R                   R                  R                  5       nUR                  U R                   R	                  U5      UU R
                  R                  U5      5      nUR                  U5        U R                  R                  X5        U R                  R                  X5        g 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   r3   (tensor_descriptor_type._flatten_ir_types    s    OO..<<>	66OO!!'*KKw'

 	

2))'7++G9r   c                l    SU R                   R                   SU R                  R                  5        S3$ )NTD_)r   mangler   r   s    r   r<   tensor_descriptor_type.mangle+   s0    DOO**+1T[[-?-?-A,B"EEr    N)returnstr)r$   List[ir.value]r%   intr?   zTuple[tensor_descriptor, int])r4   z
ir.builderr5   zList[ir.type]r?   None)
__name__
__module____qualname____firstlineno____annotations__r   r"   r3   r<   __static_attributes__r>   r   r   r   r      s0    !!F	:Fr   r   c                  l    \ rS rSr  S
S jrSS jr\S 5       r\S 5       r\S 5       r	\S 5       r
Srg	)r#   /   c                    Xl         [        R                  " U5      U l        [        R                  " U5      U l        [        X@R                  R                  U R                  R                  US9U l        g )N)r   r   r   )r&   ttgltupler'   r(   r   type)r   r&   r'   r(   r   r   s         r   __init__tensor_descriptor.__init__1   sO    ZZ&
zz'**:**//`d`l`l`q`q28:	r   c                    UR                  U R                  5        U R                  R                  U5        U R                  R                  U5        g r,   )r2   r&   r'   _flatten_irr(   )r   r$   s     r   rS   tensor_descriptor._flatten_ir9   s6    t{{#

w'  )r   c                .    U R                   R                  $ r,   )rO   r   r   s    r   r   tensor_descriptor.block_type>   s    yy###r   c                B    U R                   R                  R                  $ r,   )rO   r   r'   r   s    r   block_shapetensor_descriptor.block_shapeB   s    yy##)))r   c                B    U R                   R                  R                  $ r,   )rO   r   r-   r   s    r   dtypetensor_descriptor.dtypeF   s    yy##...r   c                .    U R                   R                  $ r,   )rO   r   r   s    r   r   tensor_descriptor.layoutJ   s    yyr   )r&   r'   r(   rO   N)r'   List[ttgl.tensor]r(   r_   r   r   r   r   )r$   rA   r?   rC   )rD   rE   rF   rG   rP   rS   propertyr   rX   r[   r   rI   r>   r   r   r#   r#   /   sd    :*:*
 $ $ * * / /    r   r#   c                    UR                  USS9nUR                  U5      nUR                  R                  U R                  XR                  UR                  UR                  5        g NF)require_i64)_convert_to_ir_values	to_tensorr4   %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<NPUWeWegmgtgt<@KKIr   c                    UR                  USS9nUR                  R                  U R                  XR                  5        g rb   )rd   r4   %create_async_tma_copy_local_to_globalr&   )rg   rh   srcrl   s       r   r   r   W   s:    ++Eu+EE;;K<N<NPUWaWabr   c                P    [        U 5      n UR                  R                  U 5        g r,   )r
   r4   create_async_tma_store_wait)pendingsrl   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_corerM   +triton.experimental.gluon.language._layoutsr   r	   r
   	triton._Cr   __all__r   r#   r   r   r   r>   r   r   <module>r~      s    " - - ! 7 7 I R
V dF F F@   @ 	I 	I 	c 	c
 	< 	<r   