xpiz ddlmZddlmZmZgdZGddeZed dZed dZedd Z ed dd d Z y))SwizzledSharedLayout)builtin_unwrap_if_constexpr)arriveinit invalidateMBarrierLayoutwaitc.eZdZdZddedeffd ZxZS)r z Layout for mbarrier synchronization in Ampere and later architectures. Args: ctas_per_cga (int): CTAs per CGA grouping. Defaults to 1. cta_split_num (int): CTA split factor. Defaults to 1. ctas_per_cga cta_split_numc :t|ddddg|g|gdgy)Nr)vec per_phase max_phaseorderr r cta_order)super__init__)selfr r __class__s {/opt/services/ai/voice_agent/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/nvidia/ampere/mbarrier.pyrzMBarrierLayout.__init__s2 #&(/c  )rr)__name__ __module__ __qualname____doc__intr __classcell__)rs@rr r s   S  S   rr 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++HOOr7sxLR D ) *  C  C = =  ]  ]  !T R  Rr