
    ^jY                     D   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ddd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                  p    e Zd ZU ded<   ded<   ded<   ded<   ddZe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    /var/www/ramen.bs-engineer-server.com/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                  }t        |      dk(  r| j                  j                  S dt        d |D              z  }| j                  j                  |z  S )Nr      c              3  @   K   | ]  }t        d  |D                yw)c              3  &   K   | ]	  }|d k7    yw)r   N ).0xs     r   	<genexpr>zBtensor_descriptor_type.nbytes_per_cta.<locals>.<genexpr>.<genexpr>   s     #:qAF#:s   N)any)r%   basiss     r   r'   z8tensor_descriptor_type.nbytes_per_cta.<locals>.<genexpr>   s     Su#:E#: :Ss   )r   
cga_layoutlenr   nbytessum)r   r*   num_cta_splitss      r   nbytes_per_ctaz%tensor_descriptor_type.nbytes_per_cta   sW    [[++
z?a??)))CS
SSS%%77r   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   r6   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 r1   )r   r2   r3   r4   r5   r   r6   appendr   _flatten_ir_typesr   )r   r7   outr8   tys        r   rE   z(tensor_descriptor_type._flatten_ir_types1   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   rK   ztensor_descriptor_type.mangle<   s3    DOO**,-Qt{{/A/A/C.DBGGr   N)returnstr)rL   int)r7   
ir.builderrL   zir.type)r=   List[ir.value]r>   rN   rL   zTuple[tensor_descriptor, int])r7   rO   rF   zList[ir.type]rL   None)__name__
__module____qualname____annotations__r   propertyr/   r6   r;   rE   rK   r$   r   r   r   r      sI    !!F 8 8
	:Hr   r   c                  p    e Zd Z	 	 d	dZd
dZe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@   rA   r   type)r   r?   r@   rA   r   r   s         r   __init__ztensor_descriptor.__init__B   sS    ZZ&
zz'**:$**//`d`l`l`q`q28:	r   c                    |j                  | j                         | j                  j                  |       | j                  j                  |       y r1   )rD   r?   r@   _flatten_irrA   )r   r=   s     r   r^   ztensor_descriptor._flatten_irJ   s6    t{{#

w'  )r   c                .    | j                   j                  S r1   )r[   r/   r   s    r   r/   z tensor_descriptor.nbytes_per_ctaO   s    yy'''r   c                .    | j                   j                  S r1   )r[   r   r   s    r   r   ztensor_descriptor.block_typeS   s    yy###r   c                B    | j                   j                  j                  S r1   )r[   r   r@   r   s    r   block_shapeztensor_descriptor.block_shapeW   s    yy##)))r   c                B    | j                   j                  j                  S r1   )r[   r   r2   r   s    r   dtypeztensor_descriptor.dtype[   s    yy##...r   c                .    | j                   j                  S r1   )r[   r   r   s    r   r   ztensor_descriptor.layout_   s    yyr   N)r@   List[ttgl.tensor]rA   rf   r   r   r   r
   )r=   rP   rL   rQ   )rR   rS   rT   r\   r^   rV   r/   r   rb   rd   r   r$   r   r   r<   r<   @   sx    :*:*
 ( ( $ $ * * / /    r   r<   c                B   t        |      d   }d}| j                  j                  rd}| j                  }|j                  dz  dk(  s
J d|        |j                  dz  }||z  }t        j                  ||      }	t        j                  d|      }
t        j                  ||      }|j                  |	|      }|j                  |
|      }| j                  j                  rdnd	}t        j                  || d
| | d| d| d|j                          
|       y )N   @      r   zunexpected sub-byte dtype 	_semanticzwith fp4_padded=True   z	 must be z"-byte aligned, i.e. a multiple of z for dtype=)listr   
fp4_paddedrd   primitive_bitwidthrY   	to_tensor__mod____eq__device_assertcodegen_name)desccoordfn_namearg_namerm   align_bytesrd   
elem_bytesalign	align_valzeroremis_zerorq   s                 r   _emit_alignment_checkr   d   s!   KOEK{{JJE##a'1,R0J5'.RR,))Q.J:%Eu	:I>>!y1DNN5I6E
--	Y-
7Cjjj3G,0KK,B,B(JG9Aj\(9[M J#WK0B0B0D/EGR[]r   c                R   |j                   j                  j                  rt        | |dd|       |j	                  |d      }|j                  |      }t        |      }|j                   j                  | j                  ||j                  |j                  |j                  |       y )Nr   innermost coordinaterl   Frequire_i64)	r7   optionsenable_iisanr   _convert_to_ir_valuesrs   r   %create_async_tma_copy_global_to_localr?   )tensor_descry   barrierresultpred	multicastrm   s          r   r   r   {   s      --k52OQg(1	3 ++Eu+EEt$D$Y/I;;r   c                    |j                   j                  j                  rt        | |dd|       |j	                  |d      }|j                   j                  | j                  ||j                         y )Nr   r   rl   Fr   )r7   r   r   r   r   %create_async_tma_copy_local_to_globalr?   )r   ry   srcrm   s       r   r   r      sg      --k52OQg(1	3++Eu+EE;;K<N<NPUWZWaWabr   c                P    t        |       } |j                  j                  |        y r1   )r   r7   create_async_tma_store_wait)pendingsrm   s     r   r   r      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 rk   rh   ri   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   r+   
ValueError
isinstancerd   rY   pointer_typer2   rr   make_scalarint32int64_unwrap_shaper[   r   r?   _str_to_padding_optionr
   rZ   r   is_intr   PADDING_OPTIONPAD_NANr7   create_make_tensor_descriptorr6   r<   )baser@   rA   rb   r   padding_optionrm   ndim	elem_sizecontig_dim_sizelast_strider&   r   base_handlepaddingr   r   rG   sr?   s                       r   make_tensor_descriptorr      s1    *.9N&{3Ku:DNN;D6MNN
7|t9TF*;CL>JKK
;48>RSVWbScRdeff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r1   )rz   rM   r{   rM   )TFN)r   N)r   zttgl.tensorr@   rf   rA   rf   rb   zList[ttgl.constexpr]r   r
   rL   r<   )
__future__r   typingr   r   r   dataclassesr   triton.language.corer   r	   (triton.experimental.gluon.language._coreexperimentalgluonlanguage_corerY   +triton.experimental.gluon.language._layoutsr
   r   r   	triton._Cr   __all__r   r<   r   r   r   r   r   r$   r   r   <module>r      s    " - - ! 6 7 7 I R
V d-HY -H -H`! 
 ! H]. 	 	$ 	c 	c 	< 	<
 	 <I
<I<I <I &	<I
 <I <I 	<Ir   