ó
    “~i?  ã                  ó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    Úq/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/experimental/gluon/language/nvidia/hopper/tma.pyÚ__str__Útensor_descriptor_type.__str__   s    € Ø# D§O¡OÐ#4°B°t·{±{°mÀ1Ð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   € Ø—O‘O×.Ñ.×<Ñ<Ó>ˆ	Ø×8Ñ8Ø�O‰O×!Ñ! 'Ó*ØØ�K‰K×Ñ˜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Ñ5°gÓF‰ˆØ×+Ñ+×9Ñ9¸'ÓJ‰ˆÜ! &°¿/¹/ÐRV×R]Ñ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„   € Ø—O‘O×.Ñ.×<Ñ<Ó>ˆ	Ø×6Ñ6Ø�O‰O×!Ñ! 'Ó*ØØ�K‰K×Ñ˜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   € Ø�D—O‘O×*Ñ*Ó,Ð-¨Q¨t¯{©{×/AÑ/AÓ/CÐ.DÀBÐ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   € àŒÜ—Z’Z Ó&ˆŒ
Ü—z’z 'Ó*ˆŒÜ*¨:Ç*Á*Ç/Á/Ð`d×`lÑ`l×`qÑ`qØ28ñ:ˆ�	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   € à�y‰y×#Ñ#Ð#r!   c                óB   • U R                   R                  R                  $ r#   )rT   r   r3   r   s    r   Úblock_shapeÚtensor_descriptor.block_shapeK   s   € à�y‰y×#Ñ#×)Ñ)Ð)r!   c                óB   • U R                   R                  R                  $ r#   )rT   r   r$   r   s    r   ÚdtypeÚtensor_descriptor.dtypeO   s   € à�y‰y×#Ñ#×.Ñ.Ð.r!   c                ó.   • U R                   R                  $ r#   )rT   r   r   s    r   r   Útensor_descriptor.layoutS   s   € à�y‰y×ÑÐ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]   € à×+Ñ+¨E¸uÐ+ÐE€EØ×Ñ˜tÓ$€DØ×Ñ×;Ñ;¸K×<NÑ<NÐPU×WeÑWeÐgm×gtÑgtØ<@¿K¹Kõ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:   € à×+Ñ+¨E¸uÐ+ÐE€EØ×Ñ×;Ñ;¸K×<NÑ<NÐPU×WaÑWaÕ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"  o¶R                  U[
        R                  5      PM$     nnU Vs/ s H6  o¶R                  [
        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['        XÏU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XÄ5      $ 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  € ô *¨.Ó9€NÜ& {Ó3€Käˆu‹:€DØ��N˜�NÜÐ;¸D¸6ÀÐMÓNÐNÜ
ˆ7ƒ|�tÓÜ˜9 T FÐ*;¼CÀ»L¸>ÐJÓKÐKÜ
ˆ;Ó˜4ÓÜÐ8¸¸Ð>RÔSVÐW^ÓS_ÐR`ÐaÓbÐbÜ�d—j‘j¤$×"3Ñ"3×4Ñ4Ð4Ð4Ø—
‘
×%Ñ%×8Ñ8¸AÑ=€IÜ×/Ò/°¸B±Ó@€OØÑ" RÓ'ÜØ`ÐapÐ`qÐqtÐu~Ðtð  @Cð  DSñ  D_ð  C`ð  `fð  gó
ð 	
ô ×+Ò+¨G°B©KÓ8€KØ�aÓÜÐHÈÈÐVÓWÐWá;@ÓAº5°a×"Ñ" 1¤d§j¡jÖ1¹5€EÐAÙX_Ó`ÒX_ÐST×$Ñ$¤T×%>Ò%>¸qÓ%AÄ4Ç:Á:ÖNÑX_€GÐ`ô ×$Ò$ [Ó1€Kä�d—i‘i¤×!2Ñ!2×3Ñ3Ð3Ð3Ü—’ §¡×!5Ñ!5°{ÓC€JØ—+‘+€Kà×.Ñ.¨~Ó>€Gä! &Ó)€FÜ�fÔ/×0Ñ0ð 4Ø3ó4Ð0ô —’˜EÓ"×'Ñ'€JÜ—:’:˜gÓ&×+Ñ+€LÜ	 
¸ÀfÓ	M€Bà‡y�y×Ñ×"Ñ"×$Ñ$¨´D·G±G×4JÑ4J×4RÑ4RÓ)RÜÐSÓTÐTØ×Ñ×<Ñ<Ø
�	‰	�)×#Ñ#Ó$ØÙ Ó!š5�aˆ�Œ™5Ñ!Ù"Ó#š7�aˆ�Œ™7Ñ#Øó€Fô ˜V U°ZÓ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!   