triton.experimental.gluon.language.nvidia.rubin.tcgen05_mma

triton.experimental.gluon.language.nvidia.rubin.tcgen05_mma(a, b, acc, *, use_acc=True, pred=True, multicast=False, mbarriers=None, mbarrier_preds=None, _semantic=None)

Emit a 5th generation TensorCore MMA instruction. acc = a * b + (acc if use_acc else 0)

Parameters:

  • a (shared_memory_descriptor or tensor_memory_descriptor) – Left hand side operand in shared or tensor memory.

  • b (shared_memory_descriptor) – Right hand side operand in shared memory.

  • acc (tensor_memory_descriptor) – Accumulator value in tensor memory (mutated).

  • use_acc (bool) – Whether to use the initial value of the accumulator. Defaults to True.

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

  • multicast (bool) – Whether tcgen05 commit should multicast across a CTA cluster. Defaults to False.

  • mbarriers (Sequence[shared_memory_descriptor], optional) – Barriers to signal when the operation is complete. If None, mma is synchronous. Defaults to None.

  • mbarrier_preds (Sequence[bool], optional) – Predicates for barriers. Defaults to None.