Function emit_marlin_q8_gemm 

pub fn emit_marlin_q8_gemm(nt: usize) -> String

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.