o i @svddlmZddlmZmZgdZGdddeZedddZedd d Zedd dZ ed ddddZ dS))SwizzledSharedLayout)builtin_unwrap_if_constexpr)arriveinit invalidateMBarrierLayoutwaitcs"eZdZdZdfdd ZZS)rz Layout for mbarrier synchronization in Ampere and later architectures. Args: cga_layout (List[List[int]]): CTA layout bases. Defaults to []. Ncs tjddddg|p gddS)Nr)vec per_phase max_phaseorder cga_layout)super__init__)selfr __class__l/home/kuhnn/.local/lib/python3.10/site-packages/triton/experimental/gluon/language/nvidia/ampere/mbarrier.pyrs zMBarrierLayout.__init__N)__name__ __module__ __qualname____doc__r __classcell__rrrrrsrNcCst|}|j|j|dS)z Initialize an mbarrier with a specified count. Args: mbarrier (shared_memory_descriptor): The barrier object to initialize. count (int): The initial count for the barrier. N)rbuildercreate_mbarrier_inithandle)mbarriercount _semanticrrrrs rcCs|j|jdS)z Invalidate an mbarrier, resetting its state. Args: mbarrier (shared_memory_descriptor): The barrier object to invalidate. N)rcreate_mbarrier_invalr)r r"rrrr srTrcCs>||}||}dd|D}|j|j|j|j|dS)a Wait until the mbarrier object completes its current phase. Args: mbarrier (shared_memory_descriptor): The barrier object to wait on. phase (int): The phase index to wait for. pred (bool): Predicate. Operation is skipped if predicate is False. Defaults to True. deps (Sequence[shared_memory_descriptor]): Dependent allocations barrier is waiting on. Used to track liveness of dependent allocations. Defaults to (). cSsg|]}|jqSr)r).0xrrr 8szwait..N) to_tensorrcreate_mbarrier_waitr)r phasepreddepsr"rrrr +s r )r*r"cCs&d}||}|j|j||jdS)a Arrive on an mbarrier, signaling that a thread has reached the barrier. Args: mbarrier (shared_memory_descriptor): The barrier object to arrive on. pred (bool): Predicate. Operation is skipped if predicate is False. Defaults to True. r N)r'rcreate_mbarrier_arriver)r r*r"r!rrrr<s rr)TrN) +triton.experimental.gluon.language._layoutsr(triton.experimental.gluon.language._corerr__all__rrrr rrrrrs