o ³´„iã@sTddlmZddlmZmZmZmZmZmZgd¢Z ed dd„ƒZ ed dd „ƒZ dS) é)Úbuiltin)Úasync_copy_global_to_sharedÚasync_copy_shared_to_globalÚ store_waitÚtensor_descriptorÚtensor_descriptor_typeÚmake_tensor_descriptor)Ú async_gatherÚ async_scatterrrrrrrTNcCs:| |¡}| |¡}|j |j|j|j|j|j|j¡dS)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Ú _semantic©rúj/home/kuhnn/.local/lib/python3.10/site-packages/triton/experimental/gluon/language/nvidia/blackwell/tma.pyr s ÿr cCs(| |¡}|j |j|j|j|j¡dS)aW 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)rrrÚsrcrrrrr *s r )TN)N) Ú(triton.experimental.gluon.language._corerÚ4triton.experimental.gluon.language.nvidia.hopper.tmarrrrrrÚ__all__r r rrrrÚs