
    bi?                    H   S SK Jr  S SKJrJrJr  S SKJr  S SKJ	r	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\  S           SS jj5       rg)    )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                  p    \ 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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/home/james-whalen/.local/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                    U R                   R                  R                  5       nUR                  U R                   R	                  U5      UU R
                  R                  U5      5      $ N)r   
element_tyis_int_signed!get_tensor_descriptor_layout_typeto_irr   _to_ir)r   builder	is_signeds      r   r(   tensor_descriptor_type._to_ir   sT    OO..<<>	88OO!!'*KKw'
 	
r!   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 r#   )r   r$   r%   r&   r'   r   r(   appendr   _flatten_ir_typesr   )r   r)   outr*   tys        r   r9   (tensor_descriptor_type._flatten_ir_types)   s    OO..<<>	66OO!!'*KKw'

 	

2))'7++G9r!   c                t    SU R                   R                  5        SU R                  R                  5        S3$ )NTD_)r   mangler   r   s    r   r@   tensor_descriptor_type.mangle4   s3    DOO**,-Qt{{/A/A/C.DBGGr!    N)returnstr)r)   
ir.builderrC   zir.type)r0   List[ir.value]r1   intrC   zTuple[tensor_descriptor, int])r)   rE   r:   zList[ir.type]rC   None)__name__
__module____qualname____firstlineno____annotations__r   r(   r.   r9   r@   __static_attributes__rB   r!   r   r   r      s5    !!F
	:H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/   8   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   )r2   ttgltupler3   r4   r   type)r   r2   r3   r4   r   r   s         r   __init__tensor_descriptor.__init__:   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#   )r8   r2   r3   _flatten_irr4   )r   r0   s     r   rX   tensor_descriptor._flatten_irB   s6    t{{#

w'  )r!   c                .    U R                   R                  $ r#   )rT   r   r   s    r   r   tensor_descriptor.block_typeG   s    yy###r!   c                B    U R                   R                  R                  $ r#   )rT   r   r3   r   s    r   block_shapetensor_descriptor.block_shapeK   s    yy##)))r!   c                B    U R                   R                  R                  $ r#   )rT   r   r$   r   s    r   dtypetensor_descriptor.dtypeO   s    yy##...r!   c                .    U R                   R                  $ r#   )rT   r   r   s    r   r   tensor_descriptor.layoutS   s    yyr!   )r2   r3   r4   rT   N)r3   List[ttgl.tensor]r4   rd   r   r   r   r
   )r0   rF   rC   rH   )rI   rJ   rK   rL   rU   rX   propertyr   r]   r`   r   rN   rB   r!   r   r/   r/   8   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_tensorr)   %create_async_tma_copy_global_to_localr2   )tensor_desccoordbarrierresultpred	_semantics         r   r   r   X   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 rg   )ri   r)   %create_async_tma_copy_local_to_globalr2   )rl   rm   srcrq   s       r   r   r   `   s:    ++Eu+EE;;K<N<NPUWaWabr!   c                P    [        U 5      n UR                  R                  U 5        g r#   )r   r)   create_async_tma_store_wait)pendingsrq   s     r   r   r   f   s     #H-H11(;r!   c           	     >   [        U5      n[        U5      n[        U5      nSUs=::  a  S::  d  O  [        SU S35      e[        U5      U:w  a  [        SU S[        U5       35      e[        U5      U:w  a  [        SU S[        U5       35      e[        U R                  [
        R                  5      (       d   eU R                  R                  R                  S	-  n[
        R                   " US
   5      n	X-  S:  a  [        SU	 SU SX-   S35      e[
        R                   " US
   5      n
U
S:w  a  [        SU
 35      eU Vs/ s H"  oR                  U[
        R                  5      PM$     nnU Vs/ s H6  oR                  [
        R                   " U5      [
        R                  5      PM8     nn[
        R                  " U5      n[        U R                  [
        R                  5      (       d   e[
        R                  " U R                  R                  U5      nU R                  nUR!                  U5      n[        U5      n[        U["        5      (       d   S5       e[
        R$                  " U5      R                  n[
        R$                  " U5      R                  n['        XUU5      nU R                  R                  R)                  5       (       a3  U[
        R*                  R,                  R.                  :X  a  [        S5      eUR0                  R3                  UR5                  UR0                  5      UU Vs/ s H  nUR                  PM     snU Vs/ s H  nUR                  PM     snU5      n[7        UXX5      $ s  snf s  snf s  snf s  snf )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
isinstancer`   rR   pointer_typer$   primitive_bitwidthmake_scalarint32int64_unwrap_shaperT   r   r2   _str_to_padding_optionr
   rS   r   is_intr   PADDING_OPTIONPAD_NANr)   create_make_tensor_descriptorr(   r/   )baser3   r4   r]   r   padding_optionrq   ndim	elem_sizecontig_dim_sizelast_stridexr   base_handlepaddingr   r   r;   sr2   s                       r   make_tensor_descriptorr   l   s6    *.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4

%%88A=I//B@O"R'`ap`qqtu~t  @C  DS  D_  C`  `f  g
 	
 ++GBK8KaHVWW;@A5a""1djj15EAX_`X_ST$$T%>%>q%A4::NX_G` $$[1Kdii!2!23333!5!5{CJ++K..~>G!&)Ff/00 4340 E"''J::g&++L	
f	MByy""$$DGG4J4J4R4R)RSTT<<
		)##$ !5a5!"#7a7#F VUZHH; B`0 	"#s   	)N8=NN N)TNr#   )zeroN)r   zttgl.tensorr3   rd   r4   rd   r]   zList[ttgl.constexpr]r   r
   rC   r/   )
__future__r   typingr   r   r   dataclassesr   triton.language.corer   r	   (triton.experimental.gluon.language._coreexperimentalgluonlanguage_corerR   +triton.experimental.gluon.language._layoutsr
   r   r   	triton._Cr   __all__r   r/   r   r   r   r   rB   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!   