Ë xõpi×ãóX—ddlmZddlmZmZmZmZmZgd¢Zedd„«Z edd„«Z y) é)Úbuiltin)Úasync_copy_global_to_sharedÚasync_copy_shared_to_globalÚ store_waitÚtensor_descriptorÚtensor_descriptor_type)Ú async_gatherÚ async_scatterrrrrrNcó—|j|«}|j|«}|jj|j|j|j|j|j|j«y)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 úy/opt/services/ai/voice_agent/venv/lib/python3.12/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ó²—|j|«}|jj|j|j|j|j«y)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Úsrcrs rr 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.tmarrrrrÚ__all__r r ©rrúr sEðÝ<÷õò €ð òJó ðJð$ ò ró ñ rr