~iG SSKJs Js Js Jr SSKJr SSKJrJ r /SQr "SS\5r \S Sj5r \S Sj5r \S SS .S j5rg) N)SwizzledSharedLayout)builtin_unwrap_if_constexpr)MBarrierLayoutinitwaitarrivec0^\rSrSrSrSU4SjjrSrU=r$)rzq Layout for mbarrier synchronization. Args: cga_layout (List[List[int]]): CTA layout bases. Defaults to []. c<>[TU]SSSS/U=(d /S9 g)Nr)vec per_phase max_phaseorder cga_layout)super__init__)selfr __class__s t/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/experimental/gluon/language/amd/gfx1250/mbarrier.pyrMBarrierLayout.__init__s$ Q!qPZP`^`aN)__name__ __module__ __qualname____firstlineno____doc__r__static_attributes__ __classcell__)rs@rrrsbbrrcf[U5nURRURU5 g)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 _semantics rrrs( ! 'E --hoouErcURU5nURRURUR5 g)aP 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_tensorr$create_lds_barrier_waitr&)r'phaser)s rrr"s2    &E --hoou||Lrr )r(r)c[U5nURRURU5n[R "U[R 5$)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. )rr$create_lds_barrier_arriver&ttgltensorint32)r'r(r)r&s rr r 2s? ! 'E    8 8% PF ;;vtzz **rr)(triton.experimental.gluon.language._core experimentalgluonlanguage_corer0+triton.experimental.gluon.language._layoutsrrr__all__rrrr rrrr:sj77LR 6 b) b  F  F  M  M 4+ +r