^jHddlZddlmcmcmZddlmZmZddl m Z ddl m Z m Z gdZGdde Zeddej d ej fd Ze dd Ze dd Ze ddZe d dddZy)N)constexpr_functionjit)SwizzledSharedLayout)builtin_unwrap_if_constexpr)allocate_mbarrierarriveinit invalidateMBarrierLayoutwaitcJeZdZdZdfd ZeeddedefdZ xZ S)r z Layout for mbarrier synchronization in Ampere and later architectures. 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/nvidia/ampere/mbarrier.pyrzMBarrierLayout.__init__s$ Q!qPZP`^`anum_ctastwo_ctactj|}tj|}|r|dzdk(sJd|dkDsJd||dz zdk(sJdg}|r|jdg|dz}tt t j |D]}|jd|zgt|S)z Create a multi-CTA mbarrier layout. Args: num_ctas (int): Number of CTAs. two_cta (bool): Whether the barrier should synchronize every other CTA rz&num_ctas must be even for two-CTA modeznum_ctas must be positiverznum_ctas must be a power of two)ttglrappendrangeintmathlog2r )rrbasesis rmultictazMBarrierLayout.multictas,,X6++G4 a<1$ N&N N$!|888|HqL)a/R1RR/  LL!  NHs499X./0 !A LL!Q$  !e$$rN)F) __name__ __module__ __qualname____doc__r staticmethodrr#boolr( __classcell__)rs@rr r s8b%3%%%rr batchtwo_ctasc,tj}|s|n|dz}tj|duxst|jt ||gn||g}tj tj|tj||}|S)z Helper function to allocate an mbarrier Args: two_ctas (bool): Whether the barrier should synchronize every other CTA rN)rr) r r static_assert isinstancevaluer#allocate_shared_memoryint64r r()r1r2r num_elemsshapebars rrr1s $}}H08h!mIu}D 5;;(DE+0=YKui>PE  % %  8D C Jrcft|}|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 rr r Es( ! '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@rBs rr r Rs++HOOrXs 11FLR Y#%)#%LT^^dnn&  C  C = =  ]  ]  !T R  Rr