Ë
    xõpi;  ã                   ó`   — d dl mZmZmZmZ ddlmZmZ g d¢Zedd„«       Z	edddd	œd
„«       Z
y)é   )ÚMBarrierLayoutÚinitÚ
invalidateÚwaité   )Ú_unwrap_if_constexprÚbuiltin)ÚarriveÚexpectr   r   r   r   TNc                 óž   — t        |«      }|j                  |«      }|j                  j                  | j                  ||j                  «       y)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Ú	_semantics       ú{/opt/services/ai/voice_agent/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/nvidia/hopper/mbarrier.pyr   r      ó>   € ô ! Ó'€EØ×Ñ˜tÓ$€DØ×Ñ×,Ñ,¨X¯_©_¸eÀTÇ[Á[ÕQó    é   )Úcountr   r   c                óž   — t        |«      }|j                  |«      }|j                  j                  | j                  ||j                  «       y)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   