rhCddlmZddlmZmZgdZGddeZed dZed dZed d Z edd Z ed d Z y))SwizzledSharedLayout)builtin_unwrap_if_constexpr)MBarrierLayoutinit invalidateexpectwaitarrivec*eZdZddedeffd ZxZS)r 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 /var/www/html/ai-insurance-compliance-backend/venv/lib/python3.12/site-packages/triton/experimental/gluon/language/nvidia/hopper/mbarrier.pyrzMBarrierLayout.__init__ s2 #&(/c  )rr)__name__ __module__ __qualname__intr __classcell__)rs@rrrs  S  S   rrNcft|}|jj|j|yN)rbuildercreate_mbarrier_inithandle)mbarriercount _semantics rrrs&  'E **8??EBrcN|jj|jyr")r#create_mbarrier_invalr%)r&r(s rrrs ++HOOr<sLR N  )   C C  = = R R  ] ] R Rr