PifddlmZddlmZmZmZmZmZgdZeddZ ed dZ dS) )builtin)async_copy_global_to_sharedasync_copy_shared_to_global store_waittensor_descriptortensor_descriptor_type) async_gather async_scatterrrrrrTNc||}||}|j|j|j|j|j|j|jdS)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_tensorbuildercreate_async_tma_gatherhandle) tensor_desc x_offsetsy_offsetbarrierresultpred _semantics /var/www/development/aibuddy-work/election-extract/venv/lib/python3.11/site-packages/triton/experimental/gluon/language/nvidia/blackwell/tma.pyr r ss   t $ $D""8,,H --k.@)BRT\Tceles.4mT[JJJJJc||}|j|j|j|j|jdS)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)rrrsrcrs rr r (sH""8,,H ..{/A9CSU]Udfifpqqqqqr)TN)N) (triton.experimental.gluon.language._corer4triton.experimental.gluon.language.nvidia.hopper.tmarrrrr__all__r r rrr s<<<<<<    JJJ J$  r r r  r r rr