Module codegen_cuda 

LOGOS → CUDA AOT codegen tier. Emits ONE self-contained CUDA program — a __global__ kernel that runs a whole LOGOS program on a single GPU thread, plus a host orchestrator — for the integer/bool scalar fragment of the language.

Why this tier tree-shakes, rather than wrapping the whole runtime

A Tenstorrent baby core is a RISC-V CPU, so crate::codegen_tt reuses the entire glibc-backed C body verbatim. A GPU thread is not a CPU: device code has no realloc/calloc/strdup/snprintf/strtod, which the C runtime leans on for every string, map, and big integer. And codegen_c emits the WHOLE ~1900-line runtime for every program, so lifting it wholesale would need device bignum + vsnprintf even for a float one-liner. So the CUDA tier TREE-SHAKES the C oracle (via crate::codegen_c::cshake): it lifts only the runtime slice a program actually references (see device_lift), __device__-annotating each unit, applying device OVERRIDES for the irreducibly-host ones (logos_fmt_f64 → a from-scratch Dragon4 shortest-float formatter, device_float), and REFUSING (typed) any reachable unit that still needs a device-absent facility. A program is accepted iff BOTH pass:

1. every op is in the scalar fragment (CUDA_SCALAR_OPS); and
2. every reachable runtime unit is device-liftable (verbatim, shimmed, or overridden).

What is verified, and when

• Pre-toolkit (no nvcc): the artifact is structurally a well-formed CUDA program, and oracle_c is byte-for-byte the C tier’s proven ≡VM body.
• On-GPU (gated on nvcc + a device): the emitted .cu compiles with nvcc and its output equals the tree-walker over the corpus (cuda_aot_lock / cuda_float_lock, skipped when the toolkit is absent — never a failure, mirroring the RISC-V/qemu gate).

Totality discipline (mirrors the RISC-V / Tenstorrent tiers)

Every refusal is CudaLowerError::Unsupported — a SOUND decline, never a miscompile. The cuda_aot_lock ratchet holds the supported-op floor and a shrink-only gap list so the fragment can only grow.

Re-exports

pub use chip::CudaChip;

pub use chip::GpuFamily;

pub use fused_row::emit_fused_row;

pub use fused_row::emit_swiglu;

pub use knobs::CachePrecision;

pub use knobs::CudaKnobs;

pub use knobs::KnobTradeoff;

pub use knobs::LoadWidth;

pub use knobs::MatmulPrecision;

pub use knobs::StatePrecision;

pub use knobs::STRICTLY_BETTER;

pub use matrix::emit_int8_gemm;

pub use matrix::lower_elementwise_to_cuda;

pub use matrix::lower_matmul_to_cuda;

pub use matrix::lower_matmul_to_cuda_with;

pub use matrix::lower_reduction_to_cuda;

pub use matrix::sample_gemm_data;

pub use matrix::CudaGemmArtifact;

pub use solbench::emit_fused_row_solution;

pub use solbench::SolSolution;

Modules

chip

The GPU chip / family knob — which NVIDIA silicon the CUDA backend targets.

forward_shim

What the older compiled forward runs whole (compile_to_cuda_forward*): its answer to the shared optimizer.

fused_row

The fused row-op CUDA emitter — the real one-block-per-row shared-memory reduction kernel.

gdn

Gated DeltaNet lowered to CUDA — the mechanism behind 48 of Qwen3.8-27B’s 64 layers.

gemm_tile

The CUDA GEMM tiling lattice — the search space our optimizing compiler sweeps to climb from the naive one-thread-per-output DP4A kernel toward the cuBLAS ceiling.

knobs

The CUDA backend’s optimization knob lattice.

lane

The compiled generation lane: a model’s decode step, written in LOGOS from its ModelSpec and compiled to one CUDA program that loads the checkpoint, runs the prompt and generates.

lane_link

The lane’s link: each op of the one IR, as the shared optimizer leaves it under the lane’s shim, to the launches the lane’s kernels run it with.

lane_shim

What the compiled generation lane runs whole: the CUDA lane’s answer to the shared optimizer.

matrix

The INT8 weight-stationary GEMM path of the CUDA tier — the workload a GPU (and an RDU) actually exists for: O(M×N) = A(M×K) · W(K×N) in int8, the shape every transformer projection and attention dot reduces to.

provenance

Where an emitted CUDA program came from, stated in its first line. Every whole-program emitter takes its header from header, so no program can claim a provenance it does not have — in particular, a hand-written lane cannot say it was compiled from a model’s .lg.

qwen38

Qwen3.8-27B authored in LOGOS.

qwen38_batched

════════════════════════════════════════════════════════════════════════════════════════════════ ★ BATCHED Qwen3.8-27B DECODE ★ — where the bandwidth wall stops being the wall.

qwen_batched

════════════════════════════════════════════════════════════════════════════════════════════════ BATCHED Qwen3-30B decode — the lane where “beat llama” is a landslide, not a 4% wall.

qwen_emit

The whole-forward orchestrator — composing a multi-stage LOGOS tensor-op program into ONE CUDA program that chains genuine PARALLEL kernels through shared device buffers.

solbench

SOL-ExecBench harness bridge — turn a recognized LOGOS fused-row op into a submittable solution.json plus a torch-binding CUDA kernel that conforms to the benchmark’s ABI.

Structs

CudaArtifact

A complete single-file CUDA program emitted from a LOGOS program.

Enums

CudaLowerError

An op / value outside the CUDA scalar fragment. The &str names what is not lifted. A refusal — the backend declines rather than emit wrong code (mirrors crate::codegen_riscv::RiscVLowerError / crate::codegen_tt::TtLowerError).

Constants

CUDA_MAKEFILE_TEMPLATE

Makefile glue: builds the single-file CUDA program with nvcc for the Ampere (sm_86, the RTX 3060 class) target. {KERNEL_FILE} is substituted with the emitted source name.

CUDA_NET_OPS

The OFFLINE NETWORKING ops the device serves via the loopback transport (codegen_c’s C_NET_RUNTIME, device-lifted). Connect/Listen/Sync are no-ops, MakePeer is identity, Send/Stream enqueue a 64-bit message into a per-address FIFO, Await dequeues it — so a single-node Send X to peer; Await … round-trips ≡ the VM’s offline NetInbox. Every op here is also in CUDA_SCALAR_OPS; excluded from the offload refusal below (they are now device-served, not host-only). Bound by cuda_aot_lock.

CUDA_SCALAR_OPS

The op names the CUDA scalar fragment lowers today — integer/bool arithmetic, comparisons, control flow, calls, globals, and Show. EXHAUSTIVELY enforced against the VM ISA by cuda_aot_lock (floor + shrink-only gaps), so this set can only grow. A scalar op whose value escapes to BigInt/Rational/Float/Text is still caught, by the device-cleanliness backstop — this list is necessary, not sufficient, for support.

KERNEL_FILE_NAME

The single canonical name of the emitted CUDA source (referenced by the Makefile).

LOGOS_DEVICE_FLOAT_FMT

The device __device__ C source for logos_fmt_f64 (and its bignum helpers). Substituted for the host body when the CUDA device-lift reaches logos_fmt_f64 (see super::device_lift). Self- contained: references only malloc/memcpy/isnan/isinf (all device-available).

Functions

codegen_program_cuda

Emit program as a single-file CUDA program. oracle_c is the self-contained C translation of the SAME program (from crate::compile::compile_to_c_stepped, which is crate::compile::compile_to_c’s C for a program with no tasks); its program-region compute is lifted onto the device. Returns CudaLowerError::Unsupported if the program leaves the scalar fragment (op-level or value-escape).

unsupported_reason

Some(reason) if program uses an op outside the CUDA scalar fragment. Host-only IO and the DOM get the shared single-device reason (crate::aot_unsupported); everything else that is not a scalar op is named directly.