ó
    +– j  ã                   ó`   • S SK Jr  S SKJrJrJrJrJrJr  / SQr	\SS j5       r
\SS j5       rg)	é    )Úbuiltin)Úasync_copy_global_to_sharedÚasync_copy_shared_to_globalÚ
store_waitÚtensor_descriptorÚtensor_descriptor_typeÚmake_tensor_descriptor)Úasync_gatherÚasync_scatterr   r   r   r   r   r	   Nc                 ó   • UR                  U5      nUR                  U5      nUR                  R                  U R                  UR                  UR                  UR                  UR                  UR                  5        g)a  
Asynchronously gather elements from global memory to shared memory using TMA.

Args:
    tensor_desc (tensor_descriptor): The tensor descriptor.
    x_offsets (tensor): 1D tensor of X offsets.
    y_offset (int): Scalar Y offset.
    barrier (shared_memory_descriptor): Barrier that will be signaled when the operation is complete.
    result (tensor_memory_descriptor): Result shared memory, must have NVMMASharedLayout.
    pred (bool): Scalar predicate. Operation is skipped if predicate is False. Defaults to True.
N)Ú	to_tensorÚbuilderÚcreate_async_tma_gatherÚhandle)Útensor_descÚ	x_offsetsÚy_offsetÚbarrierÚresultÚpredÚ	_semantics          Úˆ/root/GenerationalWealth/GenerationalWealth/venv/lib/python3.13/site-packages/triton/experimental/gluon/language/nvidia/blackwell/tma.pyr
   r
      sh   € ð ×Ñ˜tÓ$€DØ×"Ñ" 8Ó,€HØ×Ñ×-Ñ-¨k×.@Ñ.@À)×BRÑBRÐT\×TcÑTcÐel×esÑesØ.4¯m©m¸T¿[¹[õJó    c                 ó²   • UR                  U5      nUR                  R                  U R                  UR                  UR                  UR                  5        g)a;  
Asynchronously scatter elements from shared memory to global memory using TMA.

Args:
    tensor_desc (tensor_descriptor): The tensor descriptor.
    x_offsets (tensor): 1D tensor of X offsets.
    y_offset (int): Scalar Y offset.
    src (tensor_memory_descriptor): The source data, must be in NVMMASharedLayout.
N)r   r   Úcreate_async_tma_scatterr   )r   r   r   Úsrcr   s        r   r   r   *   sH   € ð ×"Ñ" 8Ó,€HØ×Ñ×.Ñ.¨{×/AÑ/AÀ9×CSÑCSÐU]×UdÑUdÐfi×fpÑfpÕqr   )TN)N)Ú(triton.experimental.gluon.language._corer   Ú4triton.experimental.gluon.language.nvidia.hopper.tmar   r   r   r   r   r	   Ú__all__r
   r   © r   r   Ú<module>r!      sE   ðÝ <÷÷ ò	€ð 	óJó 	ðJð$ 	óró 	ñrr   