
    hC                         S SK Jr  S SKJrJr  / SQr " S S\5      r\SS j5       r\SS j5       r\SS	 j5       r	\SS
 j5       r
\SS j5       rg)    )SwizzledSharedLayout)builtin_unwrap_if_constexpr)MBarrierLayoutinit
invalidateexpectwaitarrivec                   8   ^  \ rS rSrSS\S\4U 4S jjjrSrU =r$ )r      ctas_per_cgacta_split_numc           
      4   > [         TU ]  SSSS/U/U/S/S9  g )N   r   )vec	per_phase	max_phaseorderr   r   	cta_order)super__init__)selfr   r   	__class__s      s/var/www/fran/franai/venv/lib/python3.13/site-packages/triton/experimental/gluon/language/nvidia/hopper/mbarrier.pyr   MBarrierLayout.__init__	   s2    #&(/c 	 	
     )r   r   )__name__
__module____qualname____firstlineno__intr   __static_attributes____classcell__)r   s   @r   r   r      s    	
S 	
S 	
 	
r   r   Nc                 f    [        U5      nUR                  R                  U R                  U5        g N)r   buildercreate_mbarrier_inithandle)mbarriercount	_semantics      r   r   r      s&     'E**8??EBr   c                 N    UR                   R                  U R                  5        g r'   )r(   create_mbarrier_invalr*   )r+   r-   s     r   r   r      s    ++HOO<r   c                     [        U5      nUR                  U5      nUR                  R                  U R                  XR                  5        g r'   )r   	to_tensorr(   create_mbarrier_expectr*   )r+   bytespredr-   s       r   r	   r	       :     'Et$D,,X__e[[Qr   c                     UR                  U5      nUR                  U5      nU Vs/ s H  oUR                  PM     nnUR                  R                  U R                  UR                  UR                  U5        g s  snf r'   )r1   r*   r(   create_mbarrier_wait)r+   phaser4   depsr-   xs         r   r
   r
   '   sd    &Et$D"#dHHdD#**8??ELL$++W[\ $s   A;c                     [        U5      nUR                  U5      nUR                  R                  U R                  XR                  5        g r'   )r   r1   r(   create_mbarrier_arriver*   )r+   r,   r4   r-   s       r   r   r   /   r5   r   r'   )TN)Tr   N)+triton.experimental.gluon.language._layoutsr   (triton.experimental.gluon.language._corer   r   __all__r   r   r   r	   r
   r   r   r   r   <module>r@      s    L R
N
) 
 	C 	C
 	= 	= 	R 	R 	] 	] 	R 	Rr   