o iG @sddlmmmmZddlmZddlmZm Z gdZ GdddeZ edddZ edd d Z ed dd d dZdS)N)SwizzledSharedLayout)builtin_unwrap_if_constexpr)MBarrierLayoutinitwaitarrivecs"eZdZdZdfdd ZZS)rz Layout for mbarrier synchronization. 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__j/home/kuhnn/.local/lib/python3.10/site-packages/triton/experimental/gluon/language/amd/gfx1250/mbarrier.pyrs zMBarrierLayout.__init__N)__name__ __module__ __qualname____doc__r __classcell__rrrrrsrcCst|}|j|j|dS)a Initialize an mbarrier with a specified count. An mbarrier consists of an init count, a pending count and a phase. At initialization, the init count and pending count are initialized with the given 'count' and the phase is initialized to 0. Args: mbarrier (shared_memory_descriptor): The barrier object to initialize. count (int): The initial count for the barrier. Must be a positive integer. N)rbuildercreate_lds_barrier_inithandle)mbarriercount _semanticrrrrs rcCs ||}|j|j|jdS)al Wait until the mbarrier's phase differs from the provided phase value. This means that the given 'phase' has completed. Args: mbarrier (shared_memory_descriptor): The barrier object to wait on. phase (int): The phase value to compare against. The wait completes when the barrier's phase becomes different from this value. N) to_tensorrcreate_lds_barrier_waitr)rphaser!rrrr"s rr )r r!cCs&t|}|j|j|}t|tjS)a Arrive at an mbarrier with a specified count. The operation requires a `count` attribute of at least 1, and decreases the pending arrival count of the mbarrier by the specific count. If the pending count reaches zero, the phase changes (is decremented in a wraparound manner) and the pending count is reloaded with the init count value. Returns the mbarrier's phase prior to the "arrive" operation. Args: mbarrier (shared_memory_descriptor): Barrier to be signalled. count (int): Count to arrive with. Defaults to 1. Returns: prior phase (int): phase of mbarrier, prior to "arrive" operation. )rrcreate_lds_barrier_arriverttgltensorint32)rr r!rrrrr2srr)(triton.experimental.gluon.language._core experimentalgluonlanguage_corer&+triton.experimental.gluon.language._layoutsrrr__all__rrrrrrrrs