ó
    “~i×  ã                  óŒ  • S SK Jr  S SKJrJrJr  S SKJr  S SKJ	s  J
s  Js  Jr  S SKJrJr  S SKJrJr  \(       a  S SKJr  S SKJr  / S	Qr\" S
S9 " S S\R.                  5      5       r\ " S S\R2                  5      5       r\ S       SS jj5       r\ S     SS jj5       r\ S SS jj5       r\SSS jj5       rg)é    )Úannotations)ÚListÚTupleÚTYPE_CHECKING)Ú	dataclassN)ÚPaddedSharedLayoutÚSwizzledSharedLayout)ÚbuiltinÚ_unwrap_if_constexpr)Úir)Úshared_memory_descriptor)Ú
async_loadÚ
async_waitÚmake_tensor_descriptorÚtensor_descriptorÚtensor_descriptor_typeT)Úeqc                  ót   • \ rS 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)r   é   z!The type for a tensor descriptor.zttgl.block_typeÚ
block_typezttgl.tuple_typeÚ
shape_typeÚstrides_typeú)PaddedSharedLayout | SwizzledSharedLayoutÚlayoutc                ó<   • SU R                    SU R                   S3$ )Nztensor_descriptor<z, Ú>)r   r   ©Úselfs    Úo/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/experimental/gluon/language/amd/gfx1250/tdm.pyÚ__str__Útensor_descriptor_type.__str__   s    € Ø# D§O¡OÐ#4°B°t·{±{°mÀ1ÐEÐEó    c                ó¦   • X   nUS-  nU R                   R                  X5      u  pBU R                  R                  X5      u  pR[        X4XP5      nXb4$ )Né   )r   Ú_unflatten_irr   r   )r   ÚhandlesÚcursorÚhandleÚshapeÚstridesÚvalues          r   r%   Ú$tensor_descriptor_type._unflatten_ir   sT   € Ø‘ˆØ�!‰ˆØŸ™×5Ñ5°gÓF‰ˆØ×+Ñ+×9Ñ9¸'ÓJ‰ˆÜ! &°Ó?ˆØˆ}Ðr"   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   r3   Útensor_descriptor_type._to_ir$   sT   € Ø—O‘O×.Ñ.×<Ñ<Ó>ˆ	Ø×8Ñ8Ø�O‰O×!Ñ! 'Ó*ØØ�K‰K×Ñ˜wÓ'ó
ð 	
r"   c                ó°   • UR                  U R                  U5      5        U R                  R                  X5        U R                  R                  X5        g r.   )Úappendr3   r   Ú_flatten_ir_typesr   )r   r4   Úouts      r   r9   Ú(tensor_descriptor_type._flatten_ir_types,   s=   € Ø�
‰
�4—;‘;˜wÓ'Ô(Ø�‰×)Ñ)¨'Ô7Ø×Ñ×+Ñ+¨GÕ9r"   c           	     óà   • SU R                   R                  5        SU R                  R                  5        SU R                  R                  5        SU R                  R                  5        S3	$ )NÚTDÚ_)r   Úmangler   r   r   r   s    r   r?   Útensor_descriptor_type.mangle1   sb   € Ø�D—O‘O×*Ñ*Ó,Ð-¨Q¨t¯©×/EÑ/EÓ/GÐ.HÈÈ$×J[ÑJ[×JbÑJbÓJdÐIeÐefÐgk×grÑgr×gyÑgyÓg{Ðf|Ð|~ÐÐr"   © N)ÚreturnÚstr)r&   úList[ir.value]r'   ÚintrB   zTuple[tensor_descriptor, int])r4   ú
ir.builderrB   zir.type)r4   rF   r:   zList[ir.type]rB   ÚNone)Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__Ú__annotations__r    r%   r3   r9   r?   Ú__static_attributes__rA   r"   r   r   r      s8   ‡ á+àÓØÓØ!Ó!Ø5Ó5ôFôô
ô:÷
@r"   r   c                  óŒ   • \ rS rSr% SrS\S'   S\S'   S\S'   S\S	'   SS
 jr\S 5       r\S 5       r	\S 5       r
\S 5       rSrg)r   é5   z4A descriptor representing a tensor in global memory.zir.valuer(   z
ttgl.tupler)   r*   r   Útypec                ó¦   • UR                  U R                  5        U R                  R                  U5        U R                  R                  U5        g r.   )r8   r(   r)   Ú_flatten_irr*   )r   r&   s     r   rS   Útensor_descriptor._flatten_ir>   s6   € Ø�‰�t—{‘{Ô#Ø�
‰
×Ñ˜wÔ'Ø�‰× Ñ  Õ)r"   c                ó.   • U R                   R                  $ r.   )rQ   r   r   s    r   r   Útensor_descriptor.block_typeC   s   € à�y‰y×#Ñ#Ð#r"   c                óB   • U R                   R                  R                  $ r.   )rQ   r   r)   r   s    r   Úblock_shapeÚtensor_descriptor.block_shapeG   s   € à�y‰y×#Ñ#×)Ñ)Ð)r"   c                óB   • U R                   R                  R                  $ r.   )rQ   r   r/   r   s    r   ÚdtypeÚtensor_descriptor.dtypeK   s   € à�y‰y×#Ñ#×.Ñ.Ð.r"   c                ó.   • U R                   R                  $ r.   )rQ   r   r   s    r   r   Útensor_descriptor.layoutO   s   € à�y‰y×ÑÐr"   rA   N)r&   rD   rB   rG   )rH   rI   rJ   rK   rL   rM   rS   Úpropertyr   rX   r[   r   rN   rA   r"   r   r   r   5   sr   ‡ á>àÓØÓØÓØ
 Ó ô*ð
 ñ$ó ð$ð ñ*ó ð*ð ñ/ó ð/ð ñ ó ó r"   r   c                óà  • [        U5      nSUs=::  a  S::  d  O   SU S35       e[        U5      U:X  d   SU S[        U5       35       e[        U5      U:X  d   SU S[        U5       35       e[        U R                  [        R                  5      (       d   S	5       e[        U5      n[        U[        [        45      (       d   S
5       e[        U[        5      (       a  UR                  S:X  d   S5       eU R                  nUR                  USS9nUR                  USS9n	[        R                  " U5      n[        R                  " U5      n[        R                  " U R                  R                  U5      n
[        X¡R                  UR                  U5      nUR!                  S5      nUR"                  R%                  UR'                  UR"                  5      XxXœ5      n[)        XÑX+5      $ )a¥  Make a tensor descriptor object.

Args:
    base (tensor): base pointer of the tensor in global memory.
    shape (List[int]): shape of the tensor.
    strides (List[int]): strides of the tensor.
    block_shape (List[int]): block shape of the tensor.
    layout (PaddedSharedLayout | SwizzledSharedLayout): the layout of the tensor in shared memory.

Returns:
    tensor_descriptor: the created tensor descriptor object
r$   é   z Expected 1 <= ndim <= 5 but got z dimensionsz	Expected z strides but got zExpected block_shape to have z dimensions but got zExpected base to be a pointerzBExpected layout to be a PaddedSharedLayout or SwizzledSharedLayoutz3Expected max_phase to be 1 for SwizzledSharedLayoutF©Úrequire_i64TÚzero)ÚlenÚ
isinstancer[   ÚttglÚpointer_typer   r   r	   Ú	max_phaser(   Ú_convert_to_ir_valuesÚtupler   rQ   r/   r   Ú_str_to_padding_optionr4   Úcreate_make_tensor_descriptorr3   r   )Úbaser)   r*   rX   r   Ú	_semanticÚndimÚbase_handleÚshape_handlesÚstride_handlesr   rQ   Úpaddingr(   s                 r   r   r   T   sË  € ô  ˆu‹:€DØ��>˜�>ÐOÐ=¸d¸VÀ;ÐOÓOˆ>Üˆw‹<˜4ÓÐR 9¨T¨FÐ2CÄCÈÃLÀ>Ð!RÓRÐÜˆ{Ó˜tÓ#ÐmÐ'DÀTÀFÐJ^Ô_bÐcjÓ_kÐ^lÐ%mÓmÐ#Ü�d—j‘j¤$×"3Ñ"3×4Ñ4ÐUÐ6UÓUÐ4ä! &Ó)€FÜ�fÔ1Ô3GÐH×IÑIð MØLóMÐIä�&Ô.×/Ñ/Ø×Ñ 1Ó$Ð[Ð&[Ó[Ð$à—+‘+€KØ×3Ñ3°EÀuÐ3ÐM€MØ×4Ñ4°WÈ$Ð4ÐO€Nä�JŠJ�uÓ€EÜ�jŠj˜Ó!€GÜ—’ §¡×!5Ñ!5°{ÓC€JÜ! *¯j©j¸'¿,¹,ÈÓO€Dà×.Ñ.¨vÓ6€GØ×Ñ×<Ñ<¸T¿[¹[È×IZÑIZÓ=[Ð]hØ=KóV€Fô ˜V¨GÓ:Ð:r"   c                ó2  • UR                  USS9nUR                  U5      nUR                  n[        U5      nUb  UR                  O[        R
                  R                  5       nUR                  R                  U R                  XbR                  UU5        g)a  Load a block of tensor specified in tensor descriptor from global memory to shared memory asynchronously.

Args:
    src (tensor_descriptor): the source tensor descriptor.
    offsets (List[int]): the offsets from the base pointer in the tensor descriptor.
    dest (shared_memory_descriptor): the shared memory destination to store the loaded data.
    pred (bool, optional): Predicate to enable or disable the load. Defaults to True.
    mbarrier (shared_memory_descriptor, optional): The barrier object to signal "arrive" on.
Frb   N)	rj   Ú	to_tensorr(   r   rg   r   r+   r4   Ú%create_async_tdm_copy_global_to_local)	ÚsrcÚoffsetsÚdestÚpredÚmbarrierro   Úoffset_handlesÚpred_handleÚmbarrier_handles	            r   r   r   €   s   € ð ×4Ñ4°WÈ%Ð4ÐP€NØ×Ñ˜tÓ$€DØ—+‘+€KÜ# HÓ-€HØ)1Ñ)=�h—o’oÄ4Ç7Á7Ç=Á=Ã?€OØ×Ñ×;Ñ;¸C¿J¹JÈ×XcÑXcÐepØ<KõMr"   c                ó„   • UR                  USS9nUR                  R                  U R                  XBR                  5        g)aW  Store a block of tensor specified in tensor descriptor from shared memory to global memory asynchronously.

Args:
    dest (tensor_descriptor): the destination tensor descriptor.
    offsets (List[int]): the offsets from the base pointer in the tensor descriptor.
    src (shared_memory_descriptor): the shared memory source to load the data.
Frb   N)rj   r4   Ú%create_async_tdm_copy_local_to_globalr(   )rz   ry   rx   ro   r}   s        r   Úasync_storer‚   •   s9   € ð ×4Ñ4°WÈ%Ð4ÐP€NØ×Ñ×;Ñ;¸D¿K¹KÈ×YcÑYcÕdr"   c                óP   • [        U 5      n UR                  R                  U 5        g)z¢Wait for the outstanding asynchronous tensor operations to complete.

Args:
    num_outstanding (int): number of outstanding async tensor operations to wait for.
N)r   r4   Úcreate_async_tdm_wait)Únum_outstandingro   s     r   r   r   £   s"   € ô +¨?Ó;€OØ×Ñ×+Ñ+¨OÕ<r"   r.   )rn   zttgl.tensorr)   ú"List[ttgl.constexpr | ttgl.tensor]r*   r†   rX   zList[ttgl.constexpr]r   r   rB   r   )TNN)rx   r   ry   r†   rz   r   r{   Úboolr|   r   rB   rG   )rz   r   ry   r†   rx   r   rB   rG   )r   N)rB   rG   )Ú
__future__r   Útypingr   r   r   Údataclassesr   Ú(triton.experimental.gluon.language._coreÚexperimentalÚgluonÚlanguageÚ_corerg   Ú+triton.experimental.gluon.language._layoutsr   r	   r
   r   Ú	triton._Cr   r   Ú__all__Ú	base_typer   Ú
base_valuer   r   r   r‚   r   rA   r"   r   Ú<module>r•      s  ðÝ "ß -Ñ -Ý !ç 7Ö 7ß `ß RæÝÝQâ
o€ñ ˆdÑô!@˜TŸ^™^ó !@ó ð!@ðH ô ˜Ÿ™ó  ó ð ð< 	ð Y]ð(;Ø$Fð(;ØUið(;à#Lð(;àarô(;ó 	ð(;ðV 	àW[ðMØðMØ,DðMØ`dôMó 	ðMð( 	àð
eØ#'ô
eó 	ð
eð 	õ=ó 	ñ=r"   