triton.experimental.gluon.language.nvidia.rubin.mbarrier.arrive

triton.experimental.gluon.language.nvidia.rubin.mbarrier.arrive(mbarrier, *, count=1, pred=True, from_cta=None, multicast_cta=0, _semantic=None)

Arrive at an mbarrier with a specified count.

When multicast_cta is non-zero, the arrive is multicast across the cluster. Each bit set in the mask identifies a CTA ID dimension to multicast along. CTA IDs a and b belong to the same equivalence class iff a & ~multicast_cta == b & ~multicast_cta; all CTAs in a class multicast to each other. Multicast requires num_ctas > 1, 0 < multicast_cta <= num_ctas - 1, and the barrier must have the identity CGA layout [[1], [2], ...]. The default value of multicast_cta is 0 (no multicast).

Parameters:

  • mbarrier (shared_memory_descriptor) – Barrier to be signalled.

  • count (int) – Count to arrive with. Defaults to 1.

  • pred (bool) – Scalar predicate. Operation is skipped if predicate is False. Defaults to True.

  • from_cta (int, optional) – Mask of CTA-ID bits preserved when routing the arrival, in [0, num_ctas - 1]. Defaults to num_ctas - 1, which arrives from each CTA to itself; 0 routes from CTA 0 to every CTA. A non-identity mask cannot be combined with multicast.

  • multicast_cta (int) – CTA broadcast dimension bits (see above). Defaults to 0 (no multicast). Must satisfy 0 < multicast_cta <= num_ctas - 1 when non-zero.