ó
    ªêbi;  ã                   ód   • S SK JrJrJrJr  SSKJrJr  / SQr\SS j5       r	\SSSS	.S
 j5       r
g)é   )ÚMBarrierLayoutÚinitÚ
invalidateÚwaité   )Ú_unwrap_if_constexprÚbuiltin)ÚarriveÚexpectr   r   r   r   TNc                 óœ   • [        U5      nUR                  U5      nUR                  R                  U R                  XR                  5        g)aU  
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Ú	_semantics       Ús/home/james-whalen/.local/lib/python3.13/site-packages/triton/experimental/gluon/language/nvidia/hopper/mbarrier.pyr   r      ó<   € ô ! Ó'€EØ×Ñ˜tÓ$€DØ×Ñ×,Ñ,¨X¯_©_¸eÇ[Á[ÕQó    é   )Úcountr   r   c                óœ   • [        U5      nUR                  U5      nUR                  R                  U R                  XR                  5        g)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   r   r   Úcreate_mbarrier_arriver   )r   r   r   r   s       r   r
   r
      r   r   )TN)Úampere.mbarrierr   r   r   r   Ú_corer   r	   Ú__all__r   r
   © r   r   Ú<module>r       sF   ðß DÓ Dß 2â
N€ð 	óRó 	ðRð 	Ø d°dô Ró 	ñRr   