§ `ƒi×ãóf—ddlmZddlmZmZmZmZmZgd¢Zedd„¦«Z ed d„¦«Z dS) é)Úbuiltin)Úasync_copy_global_to_sharedÚasync_copy_shared_to_globalÚ store_waitÚtensor_descriptorÚtensor_descriptor_type)Ú async_gatherÚ async_scatterrrrrrTNcóÔ—| |¦«}| |¦«}|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Ú _semantics ú‹/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/triton/experimental/gluon/language/nvidia/blackwell/tma.pyr r ss€ð × Ò ˜tÑ $Ô $€DØ×"Ò" 8Ñ,Ô,€HØ Ô×-Ò-¨kÔ.@À)ÔBRÐT\ÔTcÐelÔesØ.4¬m¸T¼[ñJôJðJðJðJócó’—| |¦«}|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Úsrcrs rr r (sH€ð×"Ò" 8Ñ,Ô,€HØ Ô×.Ò.¨{Ô/AÀ9ÔCSÐU]ÔUdÐfiÔfpÑqÔqÐqÐqÐqr)TN)N) Ú(triton.experimental.gluon.language._corerÚ4triton.experimental.gluon.language.nvidia.hopper.tmarrrrrÚ__all__r r ©rrúr sÂðØ<Ð<Ð<Ð<Ð<Ð<ððððððððððððððð ð ð €ð ðJðJðJñ „ðJð$ ð rð rð rñ „ð rð rð rr