Struct CudaChip 

pub struct CudaChip {
    pub name: &'static str,
    pub family: GpuFamily,
    pub sm_arch: u32,
    pub smem_kib_per_block: usize,
    pub max_threads_per_block: usize,
    pub sms: usize,
    pub mem_bytes: usize,
    pub precisions: &'static [Precision],
}

A described CUDA target part — the hardware limits and realized numeric formats that make the fastest kernel chip-specific, plus the SM arch its kernels are compiled for. Construct with a named part (CudaChip::rtx3060, CudaChip::b200, …); the emitter reads it to pick the arch flag, the shared-memory budget, and which precisions are legal.

Fields

name: &'static str

Human-readable part name — carried into the emitted kernel’s provenance header.

family: GpuFamily

The architecture family (tensor-core generation).

sm_arch: u32

SM architecture number (80 / 86 / 89 / 90 / 100) — the sm_XX the kernel compiles for.

smem_kib_per_block: usize

Opt-in dynamic shared-memory ceiling per block, in KiB (bounds the resident row/tile a kernel stages).

max_threads_per_block: usize

Threads per block ceiling.

sms: usize

Streaming multiprocessors: the blocks a launch must offer to fill the part.

mem_bytes: usize

The part’s device memory, in bytes: what a lane plans its program’s weights, caches and arrays on, less what the driver and the program’s context take (lane::residence::CONTEXT_RESERVE).

precisions: &'static [Precision]

The numeric formats this part’s tensor / CUDA cores realize — the precision legality gate (a lowering may pick a precision only if the target chip provides it).

Implementations

impl CudaChip

pub fn rtx3060() -> Self

RTX 3060 — the local dev GPU: Ampere, sm_86, 28 SMs, ~99 KiB opt-in shared memory, 12 GiB. Where we prove correctness and relative speed before spending B200 time.

pub fn a100() -> Self

A100 — Ampere data-center part, sm_80, 108 SMs, ~164 KiB opt-in shared memory, the 80 GB part.

pub fn h100() -> Self

H100 (SXM / NVL) — Hopper, sm_90, 132 SMs, ~228 KiB opt-in shared memory, 80 GB (the SXM part), fp8.

pub fn b200() -> Self

B200 — Blackwell, sm_100, 148 SMs, ~228 KiB opt-in shared memory, 192 GB, fp8 (+ NVFP4 in Phase 5). The SOL-ExecBench target.

pub fn nvcc_arch_flag(&self) -> String

The nvcc arch flag that targets this part. Hopper/Blackwell get the architecture-specific a suffix (required to unlock wgmma/TMA on Hopper and tcgen05 on Blackwell); Ampere/Ada use the plain form.

pub fn supports(&self, p: Precision) -> bool

Whether this part’s tensor / CUDA path realizes a numeric format.

pub fn runs_numerically_on(&self, host_sm: u32) -> bool

Can a kernel built for THIS chip be numerically executed on a host GPU of SM arch host_sm? Kernels are forward-compatible: a lower-arch kernel runs on a higher-arch host, never the reverse. This routes the run-vs-compile-check split — everything sm_86-legal is proven on the local 3060; a B200 kernel can only be compile-checked locally and is numerically validated on the cloud B200.

Trait Implementations

impl Clone for CudaChip

fn clone(&self) -> CudaChip

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Debug for CudaChip

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

Formats the value using the given formatter.

impl PartialEq for CudaChip

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

impl Eq for CudaChip

impl StructuralPartialEq for CudaChip