^j ddlmcmcmcmZddlmZddlmZm Z m Z m Z m Z m Z mZgdZeddZd dZed dZy) N)builtin)async_copy_global_to_sharedasync_copy_shared_to_global store_waittensor_descriptortensor_descriptor_typemake_tensor_descriptor_emit_alignment_check) async_gather async_scatterrrrrrr cb|jjjrt||fdd||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. r y_offset _semanticN)builderoptions enable_iisanr to_tensorcreate_async_tma_gatherhandle) tensor_desc x_offsetsrbarrierresultpredrs /var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/nvidia/blackwell/tma.pyr r s  --kH<_hi   t $D""8,H --k.@.@)BRBRT\TcTceleses.4mmT[[Jc tj||}tjd|}|j||}tj|d||j||}tj|d|y)Nrrz)async_scatter y_offset cannot be negativez9async_scatter x_offsets cannot have any negative elements)ttglr__ge__ device_assert)rrrzerois_nonnegatives r_emit_scatter_nonnegative_checkr$/sq~~h)>!y 1D__TY_?N~'R^gh%%di%@N~'bnwxrc0|jjjrt||fdd|t ||||j |}|jj |j|j|j|jy)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. r rrN)rrrr r$rcreate_async_tma_scatterr)rrrsrcrs rr r :s|  --kH<*`ij' 8yQ""8,H ..{/A/A9CSCSU]UdUdfifpfpqr)TN)N)(triton.experimental.gluon.language._core experimentalgluonlanguage_corerr4triton.experimental.gluon.language.nvidia.hopper.tmarrrrrr r __all__r r$r rrr0sV77<  J J*y r rr