Module 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.

The per-op emitters (super::matrix, super::fused_row) each emit a self-contained .cu (its own main that fills one input and prints a checksum). A transformer forward is a SEQUENCE of ops (rmsnorm → matmul → softmax → …), so this module does what a single-op emitter cannot: it walks the program’s ordered tensor-op stages (each desugars to one Stmt::Splice, recognized independently via the backend-agnostic crate::tensor_recognize recognizers), then emits — for each stage — a uniquely-named __global__ kernel plus a single host driver that allocates one device buffer per tensor variable, fills the inputs deterministically, launches each stage’s kernel in dataflow order (chaining by variable name), and prints the final tensor. The kernels ARE composable — same bodies the standalone emitters use, only re-parameterized (unique name, dims as args) so many coexist in one translation unit.

This is the seam the “beat llama.cpp” path needs: the forward is authored in LOGOS tensor ops, and the compiler emits one runnable device program of chained parallel kernels — not the single-thread whole- program lowering.

Re-exports

pub use crate::tensor_recognize::qwen::recognize_moe;

pub use crate::tensor_recognize::qwen::MoeSpec;

Structs

MoeWeights

A Mixture-of-Experts FFN block‘s weights, all int8-resident (Q8_0): the router [E, d], and the E experts’ gate/up [E·dff, d] and down [E·d, dff] STACKED so expert e’s rows begin at e·dff (gate/up) or e·d (down). This is the Qwen3-30B-A3B block; the SAME struct describes any (d, dff, E, K) — the emitter indexes the experts by a base-pointer offset, so 8 experts and 128 experts share one code path.

Q8Ref

A Q8_0 weight loaded resident as int8 quants + f32 block scales (the real Qwen weights: 8 GB int8 fits a 12 GB card where 32 GB of f32 would not). Dequant W[o][t] = quants[o*in+t] * scales[o*nb + t/32] — the matmul against it becomes a dequant-in-kernel GEMV, matching codegen_cuda::qwen’s certified matvec.

Qwen30bCfg

The Qwen3-30B-A3B architecture (n_layer is overridable to keep local emit/compile fast; the per-layer shape is identical to the real model).

Enums

ChainStep

One step of a fused elementwise chain applied to the register-carried value t: a binary op against an external operand (t = t ⊕ ext[idx]), or a unary SiLU (t = silu(t), no operand).

ForwardError

Why a program is not a lowerable multi-stage tensor-op forward.

InitVal

A concrete element initializer value — integer or float.

Q8MatKind

How a Q8_0 matmul lowers to CUDA. At batch-1 decode a matmul is bound on the int8 WEIGHT bytes, not the math — so these three modes trade accuracy/latency, NOT weight bandwidth.

StageRefusal

Why the forward refuses one statement — the statement by its top-level index, and what it is.

Constants

GEN_DEVICE_KERNELS

Functions

collect_inits

Collect the concrete initializers a program writes — Set item N of V to <literal> with literal N and an integer/float value (1-based N, recorded 0-based) — so the orchestrator’s INPUT buffers hold the program’s own values and its GPU output equals the VM’s, not merely a deterministic fill. Recurses through desugar Splices.

emit_forward

Emit the composed CUDA program for a forward’s graph (forward_graph): one uniquely-named kernel per stage plus a host driver that allocates a buffer per variable, fills the inputs v[p] = (p % 7) - 3 (float chains scale by 0.5), launches each stage in order, and prints every element of the final stage’s output. Deterministic, self-contained, nvcc-compilable; its output is checked against a clean-room reference of the same graph.

emit_marlin_q8_gemm

Generate the full Qwen3-30B-A3B forward as a LOGOS .lg string, plus the manifests the emitter needs: weights (f32 embedding + RMSNorm weights), q8 (attention V/O + LM head, int8), moe (per-layer FFN, keyed by the block’s FFN output var). dir is where the weight files are (real /workspace/... on the pod, a placeholder for compile-only use). This is THE canonical authored-in-LOGOS 30B program: the pod run, the local emit lock, and every printed copy of the program come from it. The device-state kernels for autoregressive generation — a mirror of the proven 8B cuda/qwen_forward.cu decode loop, with f32 block scales (the extractor writes f32 .s files; scales are a negligible bandwidth fraction). Position + token live in device memory (d_pos/d_tok) so ONE captured graph replays every token. The GEMVs are the SAME uint4-coalesced DP4A kernels that beat llama on the single-token forward. ★ M5: the compiler EMITS the certified Marlin-class Q8_0 GEMM (it is not a hand-written .cu). Given the auto-tuned subtile count nt (the fastest golden config the auto-tuner picked for this shape/GPU — M2), this lowers the recognized matmul of a Q8_0 weight into a register-dequant tensor-core kernel: mma.sync.m16n8k32.s8 (int8 tensor core), the mma-k==32 aligned to the Q8_0 scale block (SAT-certified, M3), per-block dequant scale applied in REGISTERS in the same float grouping as our dp4a GEMV (golden-bounded to 4.2e-7, M4). This is the fusion pass in action: GEMM + per-block-dequant + scale fused into ONE kernel, and the compiler right-sizes nt per shape rather than a hand-tuned constant. Beats dp4a’s 7.5x wall on batch.

emit_moe_forward

Emit a self-contained CUDA program for ONE Mixture-of-Experts FFN block over real int8 weights: router GEMV → top-K route (softmax, keep K, renormalize) → each expert’s SwiGLU with an EXPERT-INDEXED weight base → weighted combine. The weights stream from files (int8 quants + f32 block scales), never a 4×-larger f32 buffer, so the 30B block (~30 GB int8) fits a 48 GB card. Dequant-in-kernel GEMV (f32 activation × int8 weight, f64 accumulate) matches the LOGOS moe op on the VM exactly; the whole block is captured as one CUDA graph. x is the block input (d floats, inlined). The experts are UNROLLED in the host launch list (not the kernels), so E scales without changing a kernel — the same emitter serves 8B-scale and 30B-scale.

emit_qwen30b_generate

Emit a self-contained CUDA program that AUTOREGRESSIVELY GENERATES from the real Qwen3-30B-A3B weights: the per-token math is the recognized LOGOS forward (RMSNorm → Q/K/V → per-head QK-norm + RoPE → causal attention over a KV cache → O-proj → residual → RMSNorm → 8-of-128 MoE → residual → LM head), wrapped in the compiler’s generation runtime — a prompt prefill, then a T=1 decode step captured as ONE CUDA graph and replayed per token (device-side position + token). This is the “llama.cpp, but authored in LOGOS” whole demo: it prints the generated token stream and the fair tg-style tok/s (a growing KV cache, exactly like llama-bench tg128).

forward_graph

The forward of a program as the shared table fuses it under the forward’s shim (ForwardShim): every statement the front end lifts (crate::spatial::passes::recognize_stages::walk), or the program refused whole, each refusal naming its statement. The front end decides what a statement IS (an op, a declaration, a fill, a write no op computes); kernel_refusal decides whether this lane’s kernels compute it. q8 and moe name the weights it reads in Q8_0.

forward_graph_holding

forward_graph with the optimizer’s steps held held back: the program run apart where they would fuse.

gen_qwen30b_lg

moe_dp4a_kernels

The sparse fused MoE kernel set for a coalesced int8 path: the precision-independent router + top-K pair, then the single-pass (default) or two-pass (logit-exact) quantize/gate-up-SwiGLU/down-combine trio. Public so a checked-in copy of the kernel set can be pinned byte-equal to what the compiler emits today.

recognize_attention

The single-query GQA attention block as this lane’s tuple: (q, kc, vc, out, h, kv, hd, l) — the shared qwen::AttentionSpec without its position bound (this lane attends to every cached row).