
    Pi                    8   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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                  X    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dS )tensor_descriptor_typettgl.block_type
block_typezttgl.tuple_type
shape_typestrides_typer
   layoutreturnstrc                (    d| j          d| j         dS )Nztensor_descriptor<z, >)r   r   selfs    /var/www/development/aibuddy-work/election-extract/venv/lib/python3.11/site-packages/triton/experimental/gluon/language/nvidia/hopper/tma.py__str__ztensor_descriptor_type.__str__   s    EDOEEt{EEEE    handlesList[ir.value]cursorintTuple[tensor_descriptor, int]c                    ||         }|dz  }| j                             ||          \  }}| j                            ||          \  }}t          |||| j        | j                  }||fS )N   )r   )r   _unflatten_irr   tensor_descriptorr   r   )r   r"   r$   handleshapestridesvalues          r   r)   z$tensor_descriptor_type._unflatten_ir   ss    !55gvFFv+99'6JJ!&%$/RVR]^^^f}r!   builder
ir.builderoutList[ir.type]Nonec                f   | j         j                                        }|                    | j                             |          || j                            |                    }|                    |           | j        	                    ||           | j
        	                    ||           d S N)r   
element_tyis_int_signed!get_tensor_descriptor_layout_typeto_irr   _to_irappendr   _flatten_ir_typesr   )r   r/   r1   	is_signedtys        r   r<   z(tensor_descriptor_type._flatten_ir_types!   s    O.<<>>	66O!!'**Kw''
 

 	

2))'3777++GS99999r!   c                p    d| j                                          d| j                                         dS )NTD_)r   mangler   r   s    r   rB   ztensor_descriptor_type.mangle,   s7    GDO**,,GGt{/A/A/C/CGGGGr!   N)r   r   )r"   r#   r$   r%   r   r&   )r/   r0   r1   r2   r   r3   )__name__
__module____qualname____annotations__r    r)   r<   rB    r!   r   r   r      s         !!!!F F F F   	: 	: 	: 	:H H H H H Hr!   r   c                  v    e Zd ZddZddZed             Zed             Zed             Zed             Z	dS )r*   r,   List[ttgl.tensor]r-   r   r   r   r
   c                    || _         t          j        |          | _        t          j        |          | _        t          || j        j        | j        j        |          | _        d S )N)r   r   r   )r+   ttgltupler,   r-   r   type)r   r+   r,   r-   r   r   s         r   __init__ztensor_descriptor.__init__2   sY    Z&&
z'***:$*/`d`l`q28: : :			r!   r"   r#   r   r3   c                    |                     | j                   | j                            |           | j                            |           d S r5   )r;   r+   r,   _flatten_irr-   )r   r"   s     r   rP   ztensor_descriptor._flatten_ir:   sH    t{###
w'''  )))))r!   c                    | j         j        S r5   )rM   r   r   s    r   r   ztensor_descriptor.block_type?   s    y##r!   c                $    | j         j        j        S r5   )rM   r   r,   r   s    r   block_shapeztensor_descriptor.block_shapeC   s    y#))r!   c                $    | j         j        j        S r5   )rM   r   r6   r   s    r   dtypeztensor_descriptor.dtypeG   s    y#..r!   c                    | j         j        S r5   )rM   r   r   s    r   r   ztensor_descriptor.layoutK   s    yr!   N)r,   rI   r-   rI   r   r   r   r
   )r"   r#   r   r3   )
rC   rD   rE   rN   rP   propertyr   rS   rU   r   rG   r!   r   r*   r*   0   s        : : : :* * * *
 $ $ X$ * * X* / / X/     X     r!   r*   c                    |                     |d          }|                    |          }|j                            | j        ||j        |j        |j                   d S NF)require_i64)_convert_to_ir_values	to_tensorr/   %create_async_tma_copy_global_to_localr+   )tensor_desccoordbarrierresultpred	_semantics         r   r   r   P   sp    ++Eu+EEEt$$D;;K<NPUW^Wegmgt<@KI I I I Ir!   c                    |                     |d          }|j                            | j        ||j                   d S rY   )r[   r/   %create_async_tma_copy_local_to_globalr+   )r^   r_   srcrc   s       r   r   r   X   sC    ++Eu+EEE;;K<NPUWZWabbbbbr!   c                X    t          |           } |j                            |            d S r5   )r   r/   create_async_tma_store_wait)pendingsrc   s     r   r   r   ^   s,    #H--H11(;;;;;r!   )TNr5   )
__future__r   typingr   r   r   dataclassesr   triton.language.corer   r	   (triton.experimental.gluon.language._coreexperimentalgluonlanguage_corerK   +triton.experimental.gluon.language._layoutsr
   r   r   	triton._Cr   __all__r   r*   r   r   r   rG   r!   r   <module>rv      s   " " " " " " - - - - - - - - - - ! ! ! ! ! ! 6 6 6 6 6 6 6 6 7 7 7 7 7 7 7 7 7 7 7 7 7 7 7 I I I I I I R R R R R R R R 
V
V
V dH H H H HY H H H@         
      @ 	I I I 	I 	c c c 	c
 	< < < 	< < <r!   