^j<ddlmZmZmZmZmZddlmZddlm Z m Z ddl m Z gdZ e dd Ze dd d d d Ze ddZedZy ))MBarrierLayoutallocate_mbarrierinit invalidatewait)jit)_unwrap_if_constexprbuiltin)cluster) rarriveexpectsync_cluster_initfence_init_release_clusterrrrrNTc|j|}t|}|jj|j||jy)a} Expect a specific number of bytes being copied. When they are copied, the barrier is signaled. Args: mbarrier (shared_memory_descriptor): Barrier that will be signaled when the operation is complete. bytes_per_cta (int): Expected byte count per CTA. pred (bool): Scalar predicate. Operation is skipped if predicate is False. Defaults to True. N) to_tensorr buildercreate_mbarrier_expecthandle)mbarrier bytes_per_ctapred _semantics /var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/nvidia/hopper/mbarrier.pyrrs>   t $D(7M ,,X__mT[[Y)countrrct|}|j|}|jj|j||jy)a' Arrive at an mbarrier with a specified count. Args: mbarrier (shared_memory_descriptor): Barrier to be signalled. count (int): Count to arrive with. Defaults to 1. pred (bool): Scalar predicate. Operation is skipped if predicate is False. Defaults to True. N)r rrcreate_mbarrier_arriver)rrrrs rrr"s> ! 'E   t $D ,,X__eT[[Qrc8|jjy)z Fence that makes prior mbarrier initialization visible across the CTA cluster. Needs to be called together with cluster.arrive(relaxed=True) and cluster.wait. N)r*create_fence_mbarrier_init_release_cluster)rs rrr1s@@Brclttjdtjy)zK Ensure mbarrier initialization is visible across the CTA cluster. T)relaxedN)rrrrrrrr;s   NN4  LLNr)NTN)N)ampere.mbarrierrrrrr"triton.experimental.gluon._runtimer _corer r r__all__rrrrr%rrr+sxWW22   Z  Z dd R  R C Cr