NVIDIA Rubin

async_copy

clc

Cluster Launch Control (CLC) for Blackwell (SM100+) dynamic persistent kernels.

cluster

mbarrier

tma

add2

Add two tensors using a native two-lane packed instruction.

add4

Add four-lane FP8 or FP4 operands with a Rubin packed instruction.

allocate_tensor_memory

Allocate tensor memory.

async_store

Store a tensor to shared memory asynchronously and signal an mbarrier on completion.

fence_async_shared

Order generic-proxy and asynchronous-proxy shared memory accesses.

fma2

Perform a native two-lane packed fused multiply-add.

fma4

Perform a Rubin four-lane packed FP8 or FP4 fused multiply-add.

max2

Select the maximum with a native packed half-precision instruction.

min2

Select the minimum with a native packed half-precision instruction.

mma_v2

mul2

Multiply two tensors using a native two-lane packed instruction.

mul4

Multiply four-lane FP8 or FP4 operands with a Rubin packed instruction.

sub2

Subtract two tensors using a native two-lane packed instruction.

sub4

Subtract four-lane FP8 or FP4 operands with a Rubin packed instruction.

tensor_memory_descriptor

Represents a tensor memory descriptor handle for Tensor Core Gen5 operations.

tensor_memory_descriptor_type

tcgen05_commit

This instruction causes the provided mbarrier to be arrived-on with a count of 1 when all async tcgen05 MMA and copy instructions previously issued by the thread are complete.

tcgen05_copy

Start an asynchronous copy from shared memory to tensor memory.

tcgen05_mma

Emit a 5th generation TensorCore MMA instruction.

tcgen05_mma_barrier_count

Calculate the number of CTAs that will commit the tcgen05 MMA instruction.

tcgen05_mma_scaled

Emit a 5th generation TensorCore MMA scaled instruction.

TensorMemoryLayout

Describes the layout for tensor memory in Blackwell architecture.

TensorMemoryScalesLayout

Describes the layout for tensor memory scales in Rubin architecture.