
    Wi?                    :   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e	 	 d	 	 	 	 	 	 	 	 	 	 	 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                  ^    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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    /home/sietch6/trending-topics-pipeline/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                    | j                   j                  j                         }|j                  | j                   j	                  |      || j
                  j                  |            S N)r   
element_tyis_int_signed!get_tensor_descriptor_layout_typeto_irr   _to_ir)r   builder	is_signeds      r   r&   ztensor_descriptor_type._to_ir   sT    OO..<<>	88OO!!'*KKw'
 	
r   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 r!   )r   r"   r#   r$   r%   r   r&   appendr   _flatten_ir_typesr   )r   r'   outr(   tys        r   r5   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.mangle4   s3    DOO**,-Qt{{/A/A/C.DBGGr   N)returnstr)r'   
ir.builderr<   zir.type)r-   List[ir.value]r.   intr<   zTuple[tensor_descriptor, int])r'   r>   r6   zList[ir.type]r<   None)	__name__
__module____qualname____annotations__r   r&   r+   r5   r;    r   r   r   r      s5    !!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tupler0   r1   r   type)r   r/   r0   r1   r   r   s         r   __init__ztensor_descriptor.__init__:   sS    ZZ&
zz'**:$**//`d`l`l`q`q28:	r   c                    |j                  | j                         | j                  j                  |       | j                  j                  |       y r!   )r4   r/   r0   _flatten_irr1   )r   r-   s     r   rN   ztensor_descriptor._flatten_irB   s6    t{{#

w'  )r   c                .    | j                   j                  S r!   )rK   r   r   s    r   r   ztensor_descriptor.block_typeG   s    yy###r   c                B    | j                   j                  j                  S r!   )rK   r   r0   r   s    r   block_shapeztensor_descriptor.block_shapeK   s    yy##)))r   c                B    | j                   j                  j                  S r!   )rK   r   r"   r   s    r   dtypeztensor_descriptor.dtypeO   s    yy##...r   c                .    | j                   j                  S r!   )rK   r   r   s    r   r   ztensor_descriptor.layoutS   s    yyr   N)r0   List[ttgl.tensor]r1   rU   r   r   r   r
   )r-   r?   r<   rA   )
rB   rC   rD   rL   rN   propertyr   rQ   rS   r   rF   r   r   r,   r,   8   sd    :*:*
 $ $ * * / /    r   r,   c                    |j                  |d      }|j                  |      }|j                  j                  | j                  ||j                  |j                  |j                         y 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   X   s`    ++Eu+EEt$D;;K<N<NPUW^WeWegmgtgt<@KKIr   c                    |j                  |d      }|j                  j                  | j                  ||j                         y rX   )rZ   r'   %create_async_tma_copy_local_to_globalr/   )r]   r^   srcrb   s       r   r   r   `   s=    ++Eu+EE;;K<N<NPUWZWaWabr   c                P    t        |       } |j                  j                  |        y r!   )r   r'   create_async_tma_store_wait)pendingsrb   s     r   r   r   f   s     #H-H11(;r   c           	        t        |      }t        |      }t        |      }d|cxk  rdk  sn t        d| d      t        |      |k7  rt        d| dt        |             t        |      |k7  rt        d| dt        |             t        | j                  t
        j                        sJ | j                  j                  j                  d	z  }t        j                   |d
         }	|	|z  dk  rt        d|	 d| d|	|z   d      t        j                   |d
         }
|
dk7  rt        d|
       |D cg c]"  }|j                  |t
        j                        $ }}|D cg c]5  }|j                  t        j                   |      t
        j                        7 }}t        j                  |      }t        | j                  t
        j                        sJ t        j                  | j                  j                  |      }| j                  }|j!                  |      }t        |      }t        |t"              sJ d       t        j$                  |      j                  }t        j$                  |      j                  }t'        ||||      }| j                  j                  j)                         r2|t
        j*                  j,                  j.                  k(  rt        d      |j0                  j3                  |j5                  |j0                        ||D cg c]  }|j                   c}|D cg c]  }|j                   c}|      }t7        |||||      S c c}w c c}w c c}w c c}w )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 z)Expected layout to be a NVMMASharedLayoutz8Padding option `nan` is not supported for integer blocks)r   len
ValueError
isinstancerS   rI   pointer_typer"   primitive_bitwidthmake_scalarint32int64_unwrap_shaperK   r   r/   _str_to_padding_optionr
   rJ   r   is_intr   PADDING_OPTIONPAD_NANr'   create_make_tensor_descriptorr&   r,   )baser0   r1   rQ   r   padding_optionrb   ndim	elem_sizecontig_dim_sizelast_stridexr   base_handlepaddingr   r   r7   sr/   s                       r   make_tensor_descriptorr   l   s1    *.9N&{3Ku:DNN;D6MNN
7|t9TF*;CL>JKK
;48>RSVW^S_R`abbdjj$"3"3444

%%88A=I//B@O"R'`ap`qqtu~t  @C  DS  V_  D_  C`  `f  g
 	
 ++GBK8KaHVWW;@AaY""1djj1AEAX_`STy$$T%>%>q%A4::N`G` $$[1Kdii!2!2333!5!5{CJ++K..~>G!&)Ff/0 4340 E"''J::g&++L	
Jf	MByy""$DGG4J4J4R4R)RSTT<<
		)##$ !a!"#a#F VUGZHH; B`0 	"#s   >'M$+:M)!M.:M3)TNr!   )zeroN)r|   zttgl.tensorr0   rU   r1   rU   rQ   zList[ttgl.constexpr]r   r
   r<   r,   )
__future__r   typingr   r   r   dataclassesr   triton.language.corer   r	   (triton.experimental.gluon.language._coreexperimentalgluonlanguage_corerI   +triton.experimental.gluon.language._layoutsr
   r   r   	triton._Cr   __all__r   r,   r   r   r   r   rF   r   r   <module>r      s    " - - ! 6 7 7 I R
V d%HY %H %HP 
  @ 	I 	I 	c 	c
 	< 	<
 	 <I
<I<I <I &	<I
 <I <I 	<Ir   