Wi ddlmZddlmZmZgdZGddeZed dZed dZedd Z ed dd d Z y))SwizzledSharedLayout)builtin_unwrap_if_constexpr)arriveinit invalidateMBarrierLayoutwaitc$eZdZdZdfd ZxZS)r z Layout for mbarrier synchronization in Ampere and later architectures. Args: cga_layout (List[List[int]]): CTA layout bases. Defaults to []. c8t|ddddg|xsgy)Nr)vec per_phase max_phaseorder cga_layout)super__init__)selfr __class__s /home/sietch6/trending-topics-pipeline/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/nvidia/ampere/mbarrier.pyrzMBarrierLayout.__init__s$ Q!qPZP`^`aN)__name__ __module__ __qualname____doc__r __classcell__)rs@rr r sbbrr Ncft|}|jj|j|y)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 _semantics rrrs( ! 'E **8??EBrcN|jj|jy)z Invalidate an mbarrier, resetting its state. Args: mbarrier (shared_memory_descriptor): The barrier object to invalidate. N)r create_mbarrier_invalr")r#r%s rrr s++HOOr5szLR D b) b  C  C = =  ]  ]  !T R  Rr