Enum MmaKind 

pub enum MmaKind {
    Dp4aI8,
    WmmaI8,
    HmmaBf16,
    HmmaF16,
}

The tensor-core (or CUDA-core) MMA family a GEMM lowers to, keyed on operand PRECISION. Int8 (Dp4aI8 CUDA-core, WmmaI8 IMMA) already exist; HmmaBf16/HmmaF16 add the half-precision tensor-core path. (FP8/NVFP4 join later via the mma.sync/tcgen05 PTX path on Hopper/Blackwell.)

Variants

Dp4aI8

__dp4a on the CUDA cores — the universal int8 fallback (any shape, sm_61+).

WmmaI8

IMMA int8 tensor cores, 16×16×16 s8→s32 (sm_75+, 16-aligned).

HmmaBf16

HMMA bf16 tensor cores, 16×16×16 bf16→f32 (sm_80+, 16-aligned).

HmmaF16

HMMA fp16 tensor cores, 16×16×16 f16→f32 (sm_70+, 16-aligned).

Implementations

impl MmaKind

pub fn operand_ctype(self) -> &'static str

The CUDA operand element type.

pub fn acc_ctype(self) -> &'static str

The accumulator / output element type — float for the half paths, int for int8.

pub fn header(self) -> &'static str

The extra header the operand type needs.

pub fn tag(self) -> &'static str

Short tag for kernel naming / provenance.

pub fn min_arch(self) -> u32

The minimum SM architecture the instruction needs.

pub fn is_tensor_core(self) -> bool

Whether this runs on the tensor cores (WMMA) vs the CUDA cores (DP4A).

pub fn operand_bytes(self) -> usize

Bytes per operand element — 2 for bf16/f16, 1 for int8. Sizes shared-memory footprints and the cp.async elements-per-16-byte-chunk (EPC = 16 / operand_bytes).

pub fn host_convert(self, expr: &str) -> String

Host-side conversion of a float expr into this kind’s operand type.

Trait Implementations

impl Clone for MmaKind

fn clone(&self) -> MmaKind

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Debug for MmaKind

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl PartialEq for MmaKind

fn eq(&self, other: &MmaKind) -> bool

Tests for self and other values to be equal, and is used by ==.

fn ne(&self, other: &Rhs) -> bool

Tests for !=. The default implementation is almost always sufficient, and should not be overridden without very good reason.

impl Copy for MmaKind

impl Eq for MmaKind

impl StructuralPartialEq for MmaKind