triton.experimental.gluon.language.nvidia.rubin.tcgen05_mma_scaled

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

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

Parameters:

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

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

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

  • a_scale (tensor) – Scale factor for operand A.

  • b_scale (tensor) – Scale factor for operand B.

  • a_type (str) – Type of operand A. One of {“e2m1”, “e4m3”, “e5m2”}.

  • b_type (str) – Type of operand B. One of {“e2m1”, “e4m3”, “e5m2”}.

  • 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[mbarrier], 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.