^jf ddlmcmcmcmZddlmZddlmZm Z gdZ GddeZ ed dZ ed dZ ed dd d Zy) N)SwizzledSharedLayout)builtin_unwrap_if_constexpr)MBarrierLayoutinitwaitarrivec$eZdZdZdfd ZxZS)rz Layout for mbarrier synchronization. Args: cga_layout (List[List[int]]): CGA layout bases. Defaults to []. c8t|ddddg|xsgy)Nr)vec per_phase max_phaseorder cga_layout)super__init__)selfr __class__s /var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/amd/gfx1250/mbarrier.pyrzMBarrierLayout.__init__s$ Q!qPZP`^`aN)__name__ __module__ __qualname____doc__r __classcell__)rs@rrrsbbrrcft|}|jj|j|y)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 --hoouErc|j|}|jj|j|jy)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!)r"phaser$s rrr"s2    &E --hoou||Lrr )r#r$ct|}|jj|j|}t j |tj S)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 parity (0 for even, 1 for odd) 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_arriver!ttgltensorint32)r"r#r$r!s rr r 2s? ! 'E    8 8% PF ;;vtzz **rr)(triton.experimental.gluon.language._core experimentalgluonlanguage_corer++triton.experimental.gluon.language._layoutsrrr__all__rrrr rrr6sj77LR 6 b) b  F  F  M  M 4+ +r