Piz ddlmZddlmZmZgdZGddeZeddZeddZedd Z ed dd d Z dS))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 `tddddg|g|gdgdS)Nr r)vec per_phase max_phaseorderr r cta_order)super__init__)selfr r __class__s /var/www/development/aibuddy-work/election-extract/venv/lib/python3.11/site-packages/triton/experimental/gluon/language/nvidia/ampere/mbarrier.pyrzMBarrierLayout.__init__sK #&(/c      )r r )__name__ __module__ __qualname____doc__intr __classcell__)rs@rr r sZ   S  S           rr Ncdt|}|j|j|dS)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 rrrs2 ! ' 'E **8?EBBBBBrcD|j|jdS)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#++HO<<<<zwait..As # # #AH # # #rN) to_tensorr"create_mbarrier_waitr$)r%phasepreddepsr's rr r 4sg    & &E   t $ $D # #d # # #D **8?EL$+W[\\\\\r)r3r'cd}||}|j|j||jdS)a Arrive on an mbarrier, signaling that a thread has reached the barrier. Args: mbarrier (shared_memory_descriptor): The barrier object to arrive on. pred (bool): Predicate. Operation is skipped if predicate is False. Defaults to True. r N)r0r"create_mbarrier_arriver$)r%r3r'r&s rrrEsA E   t $ $D ,,X_eT[QQQQQr)N)Tr*N) +triton.experimental.gluon.language._layoutsr(triton.experimental.gluon.language._corerr__all__r rrr rr*rrr:sLLLLLLRRRRRRRR D D D     )   *  C C C  C === =  ] ] ]  ]  !T R R R R  R R Rr