Enum SassClass 

pub enum SassClass {
    Alu,
    Fp,
    Fp64,
    Mufu,
    Tensor,
    Memory,
    AsyncCopy,
    Store,
    Control,
    Uniform,
    Other,
}

Execution-unit class of a SASS opcode — what a scheduler must know to place it and set control codes.

Variants

Alu

Fixed-latency integer/logic ALU (IMAD, IADD3, LEA, SHF, LOP3, ISETP, SEL, MOV, PRMT, …).

Fp

Fixed-latency FP32 (FADD, FFMA, FMUL, FSETP, …) and FP16x2 (HADD2, HFMA2, …).

Fp64

Fixed-latency FP64 (DADD, DFMA, DMUL, …) — longer than FP32.

Mufu

Multi-function unit (MUFU: rcp/sqrt/sin) — long fixed latency.

Tensor

Tensor core matrix-multiply-accumulate (IMMA/HMMA) — fixed throughput (~8 cyc/int8 IMMA).

Memory

Variable-latency memory that SETS A WRITE BARRIER (LDG/LDS/LDSM/LDL/LD/ATOMG/CS2R/S2R/…).

AsyncCopy

Asynchronous copy global->shared (LDGSTS = cp.async) — completion tracked by the pipeline, not a barrier.

Store

Store (STG/STS/STL/ST) — no result register.

Control

Control-flow / barriers / sync (BRA/BSSY/BSYNC/BAR/WARPSYNC/EXIT/CALL/RET/YIELD/MEMBAR/…).

Uniform

Uniform-datapath (scalar per-warp) ops (ULDC/UIADD3/UMOV/USHF/…).

Other

Anything not yet classified.

Implementations

impl SassClass

pub fn fixed_latency(&self) -> u8

Fixed issue-to-RESULT latency in cycles — the delay before a dependent consumer may read the result. CALIBRATED by microbenchmark on this sm_86 GPU (clock64 dependent chains), NOT guessed: ALU/FP ~4.2-4.5 (round to 5), int8 IMMA = 24 (MEASURED — a guessed 8 under-stalls the cross-iteration accumulate and hangs). Variable-latency classes (Memory/AsyncCopy) use write barriers instead of a fixed stall.

pub fn throughput(&self) -> u8

Functional-unit THROUGHPUT — minimum cycles between consecutive ops on the same pipe even with no register dependency. CALIBRATED on sm_86: ALU ~1.2, FP32 ~2.7, int8 IMMA ~16 for a single warp (ptxas uses ~8 assuming >=2 warps interleave the tensor pipe). Memory issue ~2. This is the constraint that serializes independent IMMAs and independent FFMAs — the register-only model misses it.

pub fn variable_result_latency(&self) -> u32

Measured VARIABLE-latency result delay in cycles for barrier-tracked ops (L1/shared hit): LDG/LDSM ~45 on sm_86 (pointer-chasing microbenchmark). The scheduler doesn’t stall for this (a barrier does), but it uses it to know these results are FAR from ready — a consumer MUST wait on the scoreboard, never a stall.

pub fn is_variable_latency(&self) -> bool

Whether this op’s result is produced with VARIABLE latency and therefore must set a write barrier (scoreboard) that dependent consumers wait on — the key fact a scheduler needs for correctness.

Trait Implementations

impl Clone for SassClass

fn clone(&self) -> SassClass

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Debug for SassClass

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

Formats the value using the given formatter.

impl Hash for SassClass

fn hash<__H: Hasher>(&self, state: &mut __H)

Feeds this value into the given Hasher.

fn hash_slice<H>(data: &[Self], state: &mut H)

where H: Hasher, Self: Sized,

Feeds a slice of this type into the given Hasher.

impl PartialEq for SassClass

fn eq(&self, other: &SassClass) -> 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 SassClass

impl Eq for SassClass

impl StructuralPartialEq for SassClass