AMD CDNA 5

async_copy

cluster

mbarrier

tdm

buffer_load

AMD buffer load from global memory via a scalar base pointer and a tensor of offsets instead of a tensor of pointers.

buffer_store

AMD buffer store a tensor directly to global memory via a scalar base pointer and a tensor of offsets instead of a tensor of pointers.

get_wmma_scale_layout

Get the scale layout for WMMA scaled operands.

load_shared_fp4_repacked

Load M/N-packed fp4 bytes from shared memory into a K-packed WMMA dot operand layout.

make_partitioned_dot_layouts

Create partitioned shared memory layouts and WMMA layout for a CDNA5 GEMM in order to avoid LDS partition conflicts.

scaled_downcast

Scale and convert FP16, BF16, or FP32 values to a low-precision MX format, dividing by the raw E8M0 scale payload (int8 or uint8).

scaled_upcast

Upcast an fp4 or fp8 tensor and fold raw E8M0 scale payload into the CDNA5 scaled-upcast op.

wmma

Computes matrix multiplication a * b + acc using an AMD WMMA instruction.

wmma_scaled

AMD Scaled WMMA operation.

PartitionedSharedLayout

Represents a partitioned shared memory layout that splits a tensor across multiple physical shared memory partitions.