
    `i;                     n    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
dS )   )MBarrierLayoutinit
invalidatewait   )_unwrap_if_constexprbuiltin)arriveexpectr   r   r   r   TNc                     t          |          }|                    |          }|j                            | j        ||j                   dS )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       /home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/triton/experimental/gluon/language/nvidia/hopper/mbarrier.pyr   r      I     !''Et$$D,,X_eT[QQQQQ       )countr   r   c                    t          |          }|                    |          }|j                            | j        ||j                   dS )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       s    D D D D D D D D D D D D 2 2 2 2 2 2 2 2
N
N
N 	R R R 	R 	dd R R R R 	R R Rr   