o ³´„i;ã@sXddlmZmZmZmZddlmZmZgd¢Zed dd„ƒZ ed ddd œd d „ƒZ dS)é)ÚMBarrierLayoutÚinitÚ invalidateÚwaité)Ú_unwrap_if_constexprÚbuiltin)ÚarriveÚexpectrrrrTNcCó*t|ƒ}| |¡}|j |j||j¡dS)am 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 (int): Expected byte count. pred (bool): Scalar predicate. Operation is skipped if predicate is False. Defaults to True. N)rÚ to_tensorÚbuilderÚcreate_mbarrier_expectÚhandle)ÚmbarrierÚbytesÚpredÚ _semantic©rúl/home/kuhnn/.local/lib/python3.10/site-packages/triton/experimental/gluon/language/nvidia/hopper/mbarrier.pyr ó r é)ÚcountrrcCr )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)rr r Úcreate_mbarrier_arriver)rrrrrrrr rr )TN) Úampere.mbarrierrrrrÚ_corerrÚ__all__r r rrrrÚs