Ë
    ŒÓ^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    € Ø# D§O¡OÐ#4°B°t·{±{°mÀ1Ð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   è ø€ Ò#:¨q A¨¥FÑ#:ùs   ‚N)Úany)r%   Úbasiss     r   r'   z8tensor_descriptor_type.nbytes_per_cta.<locals>.<genexpr>   s   è ø€ ÒS¸u¤Ñ#:°EÔ#:× :ÑSùs   ‚)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   € Ø—O‘O×.Ñ.×<Ñ<Ó>ˆ	Ø×8Ñ8Ø�O‰O×!Ñ! 'Ó*ØØ�K‰K×Ñ˜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Ñ5°g¸vÓF‰ˆˆvØ×+Ñ+×9Ñ9¸'À6ÓJ‰ˆ�Ü! &¨%°¸$¿/¹/ÐRV×R]Ñ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ˆ   € Ø—O‘O×.Ñ.×<Ñ<Ó>ˆ	Ø×6Ñ6Ø�O‰O×!Ñ! 'Ó*ØØ�K‰K×Ñ˜wÓ'ó
ˆð
 	�
‰
�2ŒØ�‰×)Ñ)¨'°3Ô7Ø×Ñ×+Ñ+¨G°SÕ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   € Ø�D—O‘O×*Ñ*Ó,Ð-¨Q¨t¯{©{×/AÑ/AÓ/CÐ.DÀBÐ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   € àˆŒÜ—Z‘Z Ó&ˆŒ
Ü—z‘z 'Ó*ˆŒÜ*¨:À$Ç*Á*Ç/Á/Ð`d×`lÑ`l×`qÑ`qØ28ô:ˆ�	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   € à�y‰y×'Ñ'Ð'r   c                ó.   — | j                   j                  S r1   )r[   r   r   s    r   r   ztensor_descriptor.block_typeS   s   € à�y‰y×#Ñ#Ð#r   c                óB   — | j                   j                  j                  S r1   )r[   r   r@   r   s    r   Úblock_shapeztensor_descriptor.block_shapeW   s   € à�y‰y×#Ñ#×)Ñ)Ð)r   c                óB   — | j                   j                  j                  S r1   )r[   r   r2   r   s    r   Údtypeztensor_descriptor.dtype[   s   € à�y‰y×#Ñ#×.Ñ.Ð.r   c                ó.   — | j                   j                  S r1   )r[   r   r   s    r   r   ztensor_descriptor.layout_   s   € à�y‰y×ÑÐ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˜‰O€EØ€KØ‡{�{×ÒØˆØ�J‰J€EØ×#Ñ# aÑ'¨1Ò,ÐRÐ0JÈ5È'Ð.RÓRÐ,Ø×)Ñ)¨QÑ.€JØ˜:Ñ%€Eä—‘˜u°	Ô:€IÜ�>‰>˜! yÔ1€Dä�N‰N˜5¨IÔ6€EØ
�-‰-˜	¨Yˆ-Ó
7€CØ�j‰j˜¨ˆjÓ3€Gà,0¯K©K×,BÒ,BÑ(È€JÜ×ÑØ�G�9˜A˜j˜\¨(¨°9¸[¸Mð JØ#˜W K°×0BÑ0BÓ0DÐ/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•   € à×Ñ× Ñ ×-Ò-Ü˜k¨5Ð2OÐQgØ(1õ	3ð ×+Ñ+¨E¸uÐ+ÓE€EØ×Ñ˜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   € à×Ñ× Ñ ×-Ò-Ü˜k¨5Ð2OÐQgØ(1õ	3à×+Ñ+¨E¸uÐ+ÓE€EØ×Ñ×;Ñ;¸K×<NÑ<NÐPUÐWZ×WaÑWaÕ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  € ô *¨.Ó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ÐWbÓScÐRdÐeÓfÐfÜ�d—j‘j¤$×"3Ñ"3Ô4Ð4Ð4Ø—
‘
×%Ñ%×8Ñ8¸AÑ=€IÜ×/Ñ/°¸B±Ó@€OØ˜Ñ" RÒ'ÜØ`ÐapÐ`qÐqtÐu~Ðtð  @Cð  DSð  V_ñ  D_ð  C`ð  `fð  gó
ð 	
ô ×+Ñ+¨G°B©KÓ8€KØ�aÒÜÐHÈÈÐVÓWÐWà;@ÖA°aˆY×"Ñ" 1¤d§j¡jÕ1ÐA€EÐAØX_Ö`ÐSTˆy×$Ñ$¤T×%>Ñ%>¸qÓ%AÄ4Ç:Á:ÕNÐ`€GÐ`ô ×$Ñ$ [Ó1€Kä�d—i‘i¤×!2Ñ!2Ô3Ð3Ð3Ü—‘ §¡×!5Ñ!5°{ÓC€JØ—+‘+€Kà×.Ñ.¨~Ó>€Gä! &Ó)€FÜ�fÔ/Ô0ð 4Ø3ó4Ð0ô —‘˜EÓ"×'Ñ'€JÜ—:‘:˜gÓ&×+Ñ+€LÜ	 
¨J¸ÀfÓ	M€Bà‡y�y×Ñ×"Ñ"Ô$¨´D·G±G×4JÑ4J×4RÑ4RÒ)RÜÐSÓTÐTØ×Ñ×<Ñ<Ø
�	‰	�)×#Ñ#Ó$ØØ Ö!�aˆ�‹Ò!Ø"Ö#�aˆ�‹Ò#Øó€Fô ˜V U¨G°ZÀÓ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   