diff --git a/.github/configs/typos-cli.toml b/.github/configs/typos-cli.toml
index db16202e..affe278b 100644
--- a/.github/configs/typos-cli.toml
+++ b/.github/configs/typos-cli.toml
@@ -29,6 +29,9 @@ extend-ignore-words-re = []
 #     rare enough that skipping a per-acronym allowlist wins.
 extend-ignore-re = [
     '`[^`\n]+`',
+    # Block-scale format names: unsigned e-m microscale exponent formats,
+    # bare or embedded in FFI constant names (VEC16_UE4M3, VEC32_UE8M0).
+    '(?i)ue\d[a-z]\d',
     '\b[A-Z][A-Z0-9]{1,4}\b',
 ]
 
diff --git a/crates/metaltile-runtime/src/device/cuda/ffi.rs b/crates/metaltile-runtime/src/device/cuda/ffi.rs
index 648ee51a..1eace3dd 100644
--- a/crates/metaltile-runtime/src/device/cuda/ffi.rs
+++ b/crates/metaltile-runtime/src/device/cuda/ffi.rs
@@ -7,9 +7,7 @@
 //! is under the NVIDIA Software License (not Apache) and the stack is alpha
 //! / Linux-only, so Phase 1 hand-rolls this small, stable surface for
 //! control. Re-evaluate `cuda-core` adoption when it leaves alpha.
-// FFI surface: bindings and constants are declared ahead of use per CUDA
-// header groupings; unused ones stay for completeness.
-#![allow(dead_code)]
+
 #![allow(non_camel_case_types)]
 
 use std::os::raw::{c_char, c_int, c_uint, c_void};
@@ -45,8 +43,14 @@ unsafe extern "C" {
     pub fn cuInit(flags: c_uint) -> CUresult;
     pub fn cuDeviceGet(device: *mut CUdevice, ordinal: c_int) -> CUresult;
     pub fn cuDeviceGetAttribute(pi: *mut c_int, attrib: c_int, dev: CUdevice) -> CUresult;
+    #[allow(dead_code)]
     pub fn cuCtxCreate_v2(pctx: *mut CUcontext, flags: c_uint, dev: CUdevice) -> CUresult;
+    pub fn cuDevicePrimaryCtxRetain(pctx: *mut CUcontext, dev: CUdevice) -> CUresult;
+    pub fn cuDevicePrimaryCtxRelease_v2(dev: CUdevice) -> CUresult;
+    pub fn cuCtxSetCurrent(ctx: CUcontext) -> CUresult;
+    #[allow(dead_code)]
     pub fn cuCtxDestroy_v2(ctx: CUcontext) -> CUresult;
+    #[allow(dead_code)]
     pub fn cuCtxSynchronize() -> CUresult;
     pub fn cuModuleLoadData(module: *mut CUmodule, image: *const c_void) -> CUresult;
     pub fn cuModuleUnload(module: CUmodule) -> CUresult;
@@ -58,6 +62,7 @@ unsafe extern "C" {
     pub fn cuFuncSetAttribute(func: CUfunction, attrib: c_int, value: c_int) -> CUresult;
     pub fn cuMemAlloc_v2(dptr: *mut CUdeviceptr, bytesize: usize) -> CUresult;
     pub fn cuMemFree_v2(dptr: CUdeviceptr) -> CUresult;
+    #[allow(dead_code)]
     pub fn cuMemcpyHtoD_v2(dst: CUdeviceptr, src: *const c_void, byte_count: usize) -> CUresult;
     pub fn cuMemcpyDtoH_v2(dst: *mut c_void, src: CUdeviceptr, byte_count: usize) -> CUresult;
     // Pinned host memory + async H2D copy — lets per-token activation uploads
@@ -71,6 +76,7 @@ unsafe extern "C" {
         byte_count: usize,
         stream: CUstream,
     ) -> CUresult;
+    pub fn cuMemsetD8Async(dst: CUdeviceptr, uc: u8, n: usize, stream: CUstream) -> CUresult;
     #[allow(clippy::too_many_arguments)]
     pub fn cuLaunchKernel(
         f: CUfunction,
@@ -122,6 +128,7 @@ unsafe extern "C" {
         flags: u64,
     ) -> CUresult;
     pub fn cuGraphLaunch(hGraphExec: CUgraphExec, hStream: CUstream) -> CUresult;
+    #[allow(dead_code)]
     pub fn cuGraphExecDestroy(hGraphExec: CUgraphExec) -> CUresult;
     pub fn cuGraphDestroy(hGraph: CUgraph) -> CUresult;
 }
@@ -157,7 +164,9 @@ pub const CUDA_R_16BF: c_int = 14;
 pub const CUBLAS_COMPUTE_32F: c_int = 68;
 
 // cublasGemmAlgo_t
+#[allow(dead_code)]
 pub const CUBLAS_GEMM_DEFAULT: c_int = -1;
+#[allow(dead_code)]
 pub const CUBLAS_GEMM_DEFAULT_TENSOR_OP: c_int = 99;
 // Explicit tensor-op algos: CUBLAS_GEMM_ALGO{N}_TENSOR_OP = 100 + N, N in 0..=15.
 // The DEFAULT (99) heuristic may pick split-K kernels that accumulate via atomics
@@ -167,6 +176,7 @@ pub const CUBLAS_GEMM_DEFAULT_TENSOR_OP: c_int = 99;
 pub const CUBLAS_GEMM_ALGO0_TENSOR_OP: c_int = 100;
 
 // cublasMath_t
+#[allow(dead_code)]
 pub const CUBLAS_DEFAULT_MATH: c_int = 0;
 pub const CUBLAS_TENSOR_OP_MATH: c_int = 1;
 
@@ -319,6 +329,23 @@ pub const CUBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET: c_int = 6;
 // cublasLtMatmulDescAttributes_t
 pub const CUBLASLT_MATMUL_DESC_TRANSA: c_int = 3;
 pub const CUBLASLT_MATMUL_DESC_TRANSB: c_int = 4;
+// Device-pointer mode + per-tensor scale pointers/modes for the fp4/fp8 GEMM
+// epilogue (cublasLtMatmulDescAttributes_t + cublasLtPointerMode_t).
+pub const CUBLASLT_MATMUL_DESC_POINTER_MODE: c_int = 2;
+pub const CUBLASLT_POINTER_MODE_DEVICE: c_int = 1;
+pub const CUBLASLT_MATMUL_DESC_A_SCALE_POINTER: c_int = 17;
+pub const CUBLASLT_MATMUL_DESC_B_SCALE_POINTER: c_int = 18;
+// D (output) scale pointer — bound for completeness of the A/B/D set; not wired
+// into a GEMM call yet.
+#[allow(dead_code)]
+pub const CUBLASLT_MATMUL_DESC_D_SCALE_POINTER: c_int = 20;
+pub const CUBLASLT_MATMUL_DESC_A_SCALE_MODE: c_int = 31;
+pub const CUBLASLT_MATMUL_DESC_B_SCALE_MODE: c_int = 32;
+// cublasLtMatmulMatrixScale_t — per-16 UE4M3 block scaling (NVFP4).
+pub const CUBLASLT_MATMUL_MATRIX_SCALE_VEC16_UE4M3: c_int = 1;
+// cudaDataType_t: E2M1 (4-bit fp4) + E4M3 (8-bit fp8).
+pub const CUDA_R_4F_E2M1: c_int = 33;
+pub const CUDA_R_8F_E4M3: c_int = 28;
 // cublasLtMatmulPreferenceAttributes_t
 pub const CUBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES: c_int = 1;
 pub const CUBLASLT_MATMUL_PREF_REDUCTION_SCHEME_MASK: c_int = 3;
@@ -440,6 +467,8 @@ unsafe extern "C" {
     ) -> nvrtcResult;
     pub fn nvrtcGetPTXSize(prog: nvrtcProgram, ptx_size_ret: *mut usize) -> nvrtcResult;
     pub fn nvrtcGetPTX(prog: nvrtcProgram, ptx: *mut c_char) -> nvrtcResult;
+    pub fn nvrtcGetCUBINSize(prog: nvrtcProgram, cubin_size_ret: *mut usize) -> nvrtcResult;
+    pub fn nvrtcGetCUBIN(prog: nvrtcProgram, cubin: *mut c_char) -> nvrtcResult;
     pub fn nvrtcGetProgramLogSize(prog: nvrtcProgram, log_size_ret: *mut usize) -> nvrtcResult;
     pub fn nvrtcGetProgramLog(prog: nvrtcProgram, log: *mut c_char) -> nvrtcResult;
     pub fn nvrtcGetErrorString(result: nvrtcResult) -> *const c_char;
diff --git a/crates/metaltile-runtime/src/device/cuda/marlin_pending.rs b/crates/metaltile-runtime/src/device/cuda/marlin_pending.rs
new file mode 100644
index 00000000..8c3aa0d0
--- /dev/null
+++ b/crates/metaltile-runtime/src/device/cuda/marlin_pending.rs
@@ -0,0 +1,75 @@
+//! Marlin-kernel runtime entry points, PENDING RECONSTRUCTION.
+//!
+//! The original implementations (device-side W4A16 Marlin batched-MoE GEMM
+//! launcher, bit-layout repack, and routing-metadata builder) lived only in
+//! an uncommitted working-tree revision that was lost to a bad sync on
+//! 2026-07-23. The Laguna engine path is unaffected: its W4A16 grouped GEMM
+//! dispatches raw kernels directly and never calls these. The NemotronH
+//! model path is the only consumer; it gets a clear error until the kernels
+//! are rebuilt (design notes survive in research/nvfp4-moe-gemm/FINDINGS.md
+//! and the ops-layer callers document the expected semantics).
+
+#![allow(clippy::too_many_arguments)]
+
+use super::{super::super::MetalTileError, CudaDevice, ffi::CUdeviceptr};
+
+fn pending(which: &str) -> MetalTileError {
+    MetalTileError::Compilation(format!(
+        "{which}: kernel pending reconstruction (see marlin_pending.rs)"
+    ))
+}
+
+impl CudaDevice {
+    pub fn moe_marlin_gemm(
+        &self,
+        _a: CUdeviceptr,
+        _b: CUdeviceptr,
+        _c: CUdeviceptr,
+        _c_tmp: CUdeviceptr,
+        _b_scales: CUdeviceptr,
+        _global_scale: CUdeviceptr,
+        _topk_weights: CUdeviceptr,
+        _sorted_token_ids: CUdeviceptr,
+        _expert_ids: CUdeviceptr,
+        _num_tokens_past_padded: CUdeviceptr,
+        _workspace: CUdeviceptr,
+        _moe_block: usize,
+        _num_experts: usize,
+        _top_k: usize,
+        _mul_topk_weights: bool,
+        _prob_m: usize,
+        _prob_n: usize,
+        _prob_k: usize,
+        _num_groups: usize,
+        _group_size: usize,
+        _sms: i32,
+        _use_fp32_reduce: bool,
+    ) -> Result<(), MetalTileError> {
+        Err(pending("moe_marlin_gemm"))
+    }
+
+    pub fn marlin_repack(
+        &self,
+        _b_q: CUdeviceptr,
+        _out: CUdeviceptr,
+        _size_k: usize,
+        _size_n: usize,
+        _sms: i32,
+        _max_shared_mem: i32,
+    ) -> Result<(), MetalTileError> {
+        Err(pending("marlin_repack"))
+    }
+
+    pub fn marlin_build_routing(
+        &self,
+        _off: CUdeviceptr,
+        _n_exp: usize,
+        _blk: usize,
+        _mt: usize,
+        _stid: CUdeviceptr,
+        _eid: CUdeviceptr,
+        _ntpp: CUdeviceptr,
+    ) -> Result<(), MetalTileError> {
+        Err(pending("marlin_build_routing"))
+    }
+}
diff --git a/crates/metaltile-runtime/src/device/cuda/mod.rs b/crates/metaltile-runtime/src/device/cuda/mod.rs
index 1e51829e..d9b04141 100644
--- a/crates/metaltile-runtime/src/device/cuda/mod.rs
+++ b/crates/metaltile-runtime/src/device/cuda/mod.rs
@@ -15,6 +15,8 @@
 //! untouched and stays the zero-config default.
 
 mod ffi;
+mod marlin_pending;
+mod nvfp4_moe;
 
 use std::{
     collections::{BTreeMap, HashMap},
@@ -165,6 +167,15 @@ pub struct CudaFunction {
     func: CUfunction,
 }
 
+/// One `cuLaunchKernel` kernel-param slot, in emitted-signature order:
+/// either a pointer param (index into `Prepared::dev_ptrs`) or a by-value
+/// scalar (index into `Prepared::scalars` — the driver copies the value
+/// from the host pointer at launch).
+enum ArgSlot {
+    Buf(usize),
+    Scalar(usize),
+}
+
 /// A compiled + resident kernel ready to launch repeatedly. See
 /// [`CudaDevice::prepare`]. Holds its device resources alive; `args()`
 /// hands out raw kernel-param pointers into `dev_ptrs`/`scalars`, so those
@@ -175,31 +186,42 @@ struct Prepared<'d> {
     dev_bufs: Vec<DeviceBuffer<'d>>,
     dev_ptrs: Vec<CUdeviceptr>,
     scalars: Vec<Vec<u8>>,
+    arg_slots: Vec<ArgSlot>,
     out_meta: Vec<Option<(String, usize)>>,
     shared_bytes: u32,
 }
 
 impl Prepared<'_> {
-    /// Build the `cuLaunchKernel` param array: param device-ptrs first, then
-    /// scalar values, each a raw pointer into our stable `dev_ptrs`/`scalars`.
+    /// Build the `cuLaunchKernel` param array in emitted-signature order,
+    /// each entry a raw pointer into our stable `dev_ptrs`/`scalars`.
     fn args(&self) -> Vec<*mut c_void> {
-        let mut args: Vec<*mut c_void> =
-            Vec::with_capacity(self.dev_ptrs.len() + self.scalars.len());
-        for p in &self.dev_ptrs {
-            args.push(p as *const CUdeviceptr as *mut c_void);
-        }
-        for s in &self.scalars {
-            args.push(s.as_ptr() as *mut c_void);
-        }
-        args
+        self.arg_slots
+            .iter()
+            .map(|slot| match slot {
+                ArgSlot::Buf(i) => &self.dev_ptrs[*i] as *const CUdeviceptr as *mut c_void,
+                ArgSlot::Scalar(i) => self.scalars[*i].as_ptr() as *mut c_void,
+            })
+            .collect()
     }
 }
 
 /// Soft cap on total bytes parked in the caching pool. Beyond this we stop
 /// retaining freed buffers (and evict on demand) so the pool can't hoard VRAM.
-/// 4 GiB is plenty of headroom for a forward's transients while leaving the
-/// 30B weights + KV resident on the 120 GB unified-memory GB10.
-const POOL_CAP_BYTES: usize = 4 * 1024 * 1024 * 1024;
+/// The 4 GiB default gives a forward's transients plenty of headroom while
+/// leaving model weights + KV resident on a unified-memory box (e.g. the
+/// 120 GB GB10); `METALTILE_POOL_CAP_MB=N` overrides it per host/workload.
+#[inline]
+fn pool_cap_bytes() -> usize {
+    use std::sync::OnceLock;
+    static CAP: OnceLock<usize> = OnceLock::new();
+    const DEFAULT: usize = 4 * 1024 * 1024 * 1024;
+    *CAP.get_or_init(|| {
+        std::env::var("METALTILE_POOL_CAP_MB")
+            .ok()
+            .and_then(|s| s.trim().parse::<usize>().ok())
+            .map_or(DEFAULT, |mb| mb * 1024 * 1024)
+    })
+}
 
 /// Round a requested allocation length up to a pool BUCKET size. A forward's
 /// transient outputs (per-layer/per-expert GEMM/attn buffers) have ever-varying
@@ -233,6 +255,7 @@ fn size_bucket(len: usize) -> usize {
 /// CUDA device + context. NVIDIA analog of `MetalDevice`.
 pub struct CudaDevice {
     ctx: CUcontext,
+    dev: CUdevice,
     cc_major: i32,
     cc_minor: i32,
     /// Caching device allocator (PyTorch-style). `cuMemAlloc`/`cuMemFree` are
@@ -248,13 +271,15 @@ pub struct CudaDevice {
     /// finished consuming it.
     ///
     /// Key = bucket size in bytes (NOT the requested len). `pooled_bytes`
-    /// tracks the total parked so we can cap the pool ([`POOL_CAP_BYTES`]) and
-    /// evict beyond it rather than hoard VRAM. Gated by `METALTILE_POOL_ALLOC`
-    /// ([`pool_enabled`]): off ⇒ direct `cuMemAlloc`/`cuMemFree` (clean A/B).
+    /// tracks the total parked so we can cap the pool ([`pool_cap_bytes`]) and
+    /// evict beyond it rather than hoard VRAM. Default-on; set
+    /// `METALTILE_POOL_ALLOC_OFF=1` ([`pool_enabled`]) for direct
+    /// `cuMemAlloc`/`cuMemFree` (clean A/B).
     pool: Mutex<HashMap<usize, Vec<CUdeviceptr>>>,
     /// Total bytes currently parked in `pool` (sum of bucket_size × count).
     pooled_bytes: Mutex<usize>,
-    /// `METALTILE_POOL_ALLOC` is set ⇒ caching allocator active. Cached once.
+    /// Caching allocator active unless `METALTILE_POOL_ALLOC_OFF` is set.
+    /// Cached once at create().
     pool_enabled: bool,
     /// Pinned host-staging buffers for async H2D: free-list by size, plus the
     /// in-flight list whose copies haven't been synced yet (reclaimed on the next
@@ -282,22 +307,50 @@ pub struct CudaDevice {
     /// heuristic / atomics-mode could not on sm_121. `(handle_ptr, workspace_ptr)`
     /// as `usize` to stay plain-Send; workspace is a single 32 MiB slab.
     cublaslt: Mutex<(usize, usize)>,
+    /// Lazily-allocated device f32 scalar `0.0`, used as the `beta` pointer for
+    /// the cuBLASLt fp4/fp8 GEMM epilogue (device pointer mode). `usize` to stay
+    /// plain-Send; see `lt_beta_zero_ptr` (in `nvfp4_moe`).
+    lt_beta_zero: Mutex<usize>,
 }
 
 /// Fixed cublasLt workspace size (32 MiB). Large enough for the heuristic to
 /// pick efficient non-split-K tensor-op kernels for the Nemotron prefill GEMMs.
 const CUBLASLT_WORKSPACE_BYTES: usize = 32 * 1024 * 1024;
 
-// The context is current on this struct's lifetime; we keep it single-
-// device, single-context (Phase 1). Send is sound because we never share
-// the raw pointers across threads concurrently in the smoke path. Sync is
-// sound for serialized submission (a higher layer holds the device in an
-// `Arc` to keep its context alive for persistent buffers; GPU work is
-// submitted from one logical owner at a time).
+// SAFETY / THREADING CONTRACT: `cuCtxCreate` makes the context current on
+// the CREATING thread only, so every driver call in this module assumes it
+// runs on the thread that called [`CudaDevice::create`]. Calls from any
+// other thread see no current context and fail with a clean
+// CUDA_ERROR_INVALID_CONTEXT (interior state is Mutex-guarded, so there is
+// no data race — just an error). Send/Sync are asserted so a higher layer
+// can hold the device in an `Arc`/static to keep the context (and thus
+// persistent device buffers) alive; all GPU submission must stay on the
+// The device holds the refcounted PRIMARY context; `ensure_current` binds it
+// to the calling thread (cuCtxSetCurrent, cached per thread) at every public
+// entry point, so cross-thread use is sound.
 unsafe impl Send for CudaDevice {}
 unsafe impl Sync for CudaDevice {}
 
+thread_local! {
+    /// The ctx this thread last bound via `ensure_current` — skip the driver
+    /// call when it's already current (the per-call cost would otherwise be
+    /// paid ~400×/token on the decode path).
+    static CURRENT_CTX: std::cell::Cell<usize> = const { std::cell::Cell::new(0) };
+}
+
 impl CudaDevice {
+    /// Bind this device's primary context to the calling thread (no-op when
+    /// already bound). Must run at the top of every public entry point that
+    /// touches the driver.
+    fn ensure_current(&self) {
+        CURRENT_CTX.with(|c| {
+            if c.get() != self.ctx as usize {
+                unsafe { cuCtxSetCurrent(self.ctx) };
+                c.set(self.ctx as usize);
+            }
+        });
+    }
+
     /// Initialize CUDA, grab device 0, create a context. Returns `Ok(None)`
     /// if no CUDA device is present (mirrors `MetalDevice::create`).
     pub fn create() -> Result<Option<Self>, MetalTileError> {
@@ -319,16 +372,28 @@ impl CudaDevice {
                 cuDeviceGetAttribute(&mut minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, dev),
                 "cuDeviceGetAttribute(cc_minor)",
             )?;
+            // Primary context (shared, refcounted) instead of cuCtxCreate:
+            // a created context is current ONLY on the creating thread,
+            // which made the Send/Sync impls a documented lie. The primary
+            // ctx + a per-thread cuCtxSetCurrent guard (ensure_current)
+            // makes cross-thread use sound.
             let mut ctx: CUcontext = ptr::null_mut();
-            cu_check(cuCtxCreate_v2(&mut ctx, 0, dev), "cuCtxCreate")?;
+            cu_check(cuDevicePrimaryCtxRetain(&mut ctx, dev), "cuDevicePrimaryCtxRetain")?;
+            cu_check(cuCtxSetCurrent(ctx), "cuCtxSetCurrent")?;
             let mut stream: CUstream = ptr::null_mut();
-            cu_check(cuStreamCreate(&mut stream, CU_STREAM_NON_BLOCKING), "cuStreamCreate")?;
+            let sres = cuStreamCreate(&mut stream, CU_STREAM_NON_BLOCKING);
+            if sres != CUDA_SUCCESS {
+                cuDevicePrimaryCtxRelease_v2(dev);
+                cu_check(sres, "cuStreamCreate")?;
+            }
             // Caching allocator DEFAULT-ON (alloc()/Drop route through it): nsys
             // showed raw cuMemAlloc/cuMemFree at 62% of CUDA API time. METALTILE_POOL_ALLOC_OFF=1
-            // restores direct driver alloc/free (clean A/B). 4GB pool cap (POOL_CAP_BYTES).
+            // restores direct driver alloc/free (clean A/B). Pool cap defaults
+            // to 4GB; METALTILE_POOL_CAP_MB overrides (pool_cap_bytes()).
             let pool_enabled = std::env::var("METALTILE_POOL_ALLOC_OFF").is_err();
             Ok(Some(CudaDevice {
                 ctx,
+                dev,
                 cc_major: major,
                 cc_minor: minor,
                 pool: Mutex::new(HashMap::new()),
@@ -340,6 +405,7 @@ impl CudaDevice {
                 capturing: std::sync::atomic::AtomicBool::new(false),
                 cublas: Mutex::new(0),
                 cublaslt: Mutex::new((0, 0)),
+                lt_beta_zero: Mutex::new(0),
             }))
         }
     }
@@ -432,6 +498,7 @@ impl CudaDevice {
         k: usize,
         dtype: metaltile_core::DType,
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         // DETERMINISTIC by default: route through cublasLt with split-K reductions
         // forbidden (REDUCTION_SCHEME_NONE). The legacy cublasGemmEx heuristic
         // picks split-K atomic-accumulate kernels for some MoE-prefill shapes →
@@ -509,7 +576,7 @@ impl CudaDevice {
     /// (HOST slices) describe each contiguous token group + its weight slab.
     /// `out[t,n] = Σ_k a[t,k]·w[eid][n,k]`. Beats cuBLAS per-expert on the skinny
     /// MoE shape. Errors if the runtime was built without CUTLASS.
-    #[allow(clippy::too_many_arguments)]
+    #[allow(clippy::too_many_arguments)] // mirrors the C entry point's signature
     pub fn moe_grouped_cutlass(
         &self,
         a: CUdeviceptr,
@@ -520,6 +587,7 @@ impl CudaDevice {
         n: usize,
         k: usize,
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         #[cfg(have_cutlass)]
         {
             unsafe extern "C" {
@@ -586,6 +654,7 @@ impl CudaDevice {
         k: usize,
         ab_dtype: metaltile_core::DType,
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         self.gemm_cublaslt(x, w, out, m, n, k, ab_dtype, metaltile_core::DType::F32)
     }
 
@@ -611,6 +680,7 @@ impl CudaDevice {
         dtype: metaltile_core::DType,
         out_dtype: metaltile_core::DType,
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         let (lt, workspace) = self.cublaslt_ctx()?;
         let cdt = match dtype {
             metaltile_core::DType::F16 => CUDA_R_16F,
@@ -762,6 +832,7 @@ impl CudaDevice {
         batch_count: usize,
         dtype: metaltile_core::DType,
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         // DETERMINISTIC by default via cublasLt (split-K reductions forbidden);
         // see gemm_cublas. METALTILE_GEMM_NONDET=1 → legacy nondeterministic path.
         if std::env::var("METALTILE_GEMM_NONDET").ok().as_deref() != Some("1") {
@@ -845,6 +916,7 @@ impl CudaDevice {
         batch_count: usize,
         dtype: metaltile_core::DType,
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         let (lt, workspace) = self.cublaslt_ctx()?;
         let cdt = match dtype {
             metaltile_core::DType::F16 => CUDA_R_16F,
@@ -1003,6 +1075,7 @@ impl CudaDevice {
         k: usize,            // shared input dim
         dtype: metaltile_core::DType,
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         let group_count = x_ptrs.len();
         assert_eq!(w_ptrs.len(), group_count);
         assert_eq!(out_ptrs.len(), group_count);
@@ -1026,9 +1099,11 @@ impl CudaDevice {
         // Each "group" has group_size[i]=1 (one GEMM with its own m).
         let n_i = n as i32;
         let k_i = k as i32;
-        // alpha/beta: single scalar shared across all groups (same as GemmEx)
-        let alpha_scalar: f32 = 1.0;
-        let beta_scalar: f32 = 0.0;
+        // alpha/beta: the grouped API takes HOST ARRAYS with one entry per
+        // group (unlike GemmEx's single scalar) — a lone scalar ref would be
+        // read group_count entries deep (OOB host read for group_count > 1).
+        let alphas: Vec<f32> = vec![1.0f32; group_count];
+        let betas: Vec<f32> = vec![0.0f32; group_count];
         let transas: Vec<c_int> = vec![CUBLAS_OP_T; group_count];
         let transbs: Vec<c_int> = vec![CUBLAS_OP_N; group_count];
         // m_array[i] = n (cm rows = output cols), n_array[i] = m_per_group[i] (cm cols = token rows)
@@ -1048,8 +1123,8 @@ impl CudaDevice {
         // cuMemAlloc + 3× cuMemcpyHtoD + 3× cuMemFree per call — each driver
         // alloc/free is a device-wide sync that serialized the stream per grouped
         // GEMM). `alloc_raw`/`free_raw_pooled` recycle the buffer across calls when
-        // METALTILE_POOL_ALLOC is set (the raw driver alloc here otherwise bypassed
-        // the caching allocator entirely).
+        // the pool is enabled (default; the raw driver alloc here otherwise
+        // bypassed the caching allocator entirely).
         let ptr_bytes = group_count * 8; // each pointer is 8 bytes (u64)
         let triple_bytes = ptr_bytes * 3;
         // Layout: [A(W) | B(X) | C(Out)] contiguous.
@@ -1092,14 +1167,14 @@ impl CudaDevice {
                 m_arr.as_ptr(),
                 n_arr.as_ptr(),
                 k_arr.as_ptr(),
-                &alpha_scalar as *const f32 as *const c_void,
+                alphas.as_ptr() as *const c_void,
                 a_dev as *const *const c_void,
                 cdt,
                 lda.as_ptr(),
                 b_dev as *const *const c_void,
                 cdt,
                 ldb.as_ptr(),
-                &beta_scalar as *const f32 as *const c_void,
+                betas.as_ptr() as *const c_void,
                 c_dev as *mut *mut c_void,
                 cdt,
                 ldc.as_ptr(),
@@ -1141,6 +1216,7 @@ impl CudaDevice {
         k: usize,
         dtype: metaltile_core::DType,
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         let batch_count = x_ptrs.len();
         assert_eq!(w_ptrs.len(), batch_count);
         assert_eq!(out_ptrs.len(), batch_count);
@@ -1231,6 +1307,7 @@ impl CudaDevice {
     /// driver JIT. Targets the device's own virtual arch
     /// (`compute_<major><minor>`); the driver JITs PTX to the live GPU.
     pub fn compile(&self, src: &str, prog_name: &str) -> Result<CudaModule, MetalTileError> {
+        self.ensure_current();
         let csrc = CString::new(src).map_err(|e| MetalTileError::Compilation(e.to_string()))?;
         let cname =
             CString::new(prog_name).map_err(|e| MetalTileError::Compilation(e.to_string()))?;
@@ -1250,10 +1327,19 @@ impl CudaDevice {
             "nvrtcCreateProgram",
         )?;
 
-        // Compile to the device's virtual architecture.
-        let arch =
+        // PTX + driver JIT is the DEFAULT: an A/B on GB10 measured it
+        // slightly ahead of native SASS for the latency-bound decode
+        // kernels (36.3 vs 35.4 tok/s) and equal on prefill. MT_NVRTC_NATIVE=1
+        // opts in to native-arch SASS via nvrtcGetCUBIN (faster first
+        // dispatch, no JIT).
+        let use_ptx = std::env::var("MT_NVRTC_NATIVE").is_err();
+        let arch = if use_ptx {
             CString::new(format!("--gpu-architecture=compute_{}{}", self.cc_major, self.cc_minor))
-                .unwrap();
+                .unwrap()
+        } else {
+            CString::new(format!("--gpu-architecture=sm_{}{}a", self.cc_major, self.cc_minor))
+                .unwrap()
+        };
         // NVRTC does not auto-include the toolkit headers (cuda_fp16.h,
         // cuda_bf16.h) — point it at <toolkit>/include.
         let cuda_root = std::env::var("CUDA_PATH")
@@ -1269,21 +1355,42 @@ impl CudaDevice {
         let cccl_path = {
             let fixed1 = format!("{cuda_root}/targets/sbsa-linux/include/cccl");
             let fixed2 = format!("{cuda_root}/targets/x86_64-linux/include/cccl");
-            // Also try the versioned cuda-13.x path (some distros install both).
-            let by_ver = std::fs::read_dir(format!(
-                "{cuda_root}-{}.{}/targets",
-                self.cc_major, self.cc_minor
-            ))
-            .ok()
-            .and_then(|mut rd| rd.next())
-            .and_then(|e| e.ok())
-            .map(|e| format!("{}/include/cccl", e.path().display()));
+            // Versioned TOOLKIT installs sit as `<root>-<toolkit-version>`
+            // siblings of the unversioned root (e.g. /usr/local/cuda-13.0);
+            // probe those, newest first. (The toolkit version is unrelated to
+            // the device's compute capability — don't derive paths from it.)
+            let by_toolkit = || -> Option<String> {
+                let root = std::path::Path::new(&cuda_root);
+                let stem = root.file_name()?.to_str()?.to_string();
+                let mut versioned: Vec<std::path::PathBuf> = std::fs::read_dir(root.parent()?)
+                    .ok()?
+                    .flatten()
+                    .map(|e| e.path())
+                    .filter(|p| {
+                        p.file_name()
+                            .and_then(|n| n.to_str())
+                            .is_some_and(|n| n.starts_with(&format!("{stem}-")))
+                    })
+                    .collect();
+                versioned.sort();
+                for dir in versioned.into_iter().rev() {
+                    if let Ok(rd) = std::fs::read_dir(dir.join("targets")) {
+                        for t in rd.flatten() {
+                            let cccl = t.path().join("include/cccl");
+                            if cccl.is_dir() {
+                                return Some(cccl.display().to_string());
+                            }
+                        }
+                    }
+                }
+                None
+            };
             if std::path::Path::new(&fixed1).exists() {
                 Some(fixed1)
             } else if std::path::Path::new(&fixed2).exists() {
                 Some(fixed2)
             } else {
-                by_ver.filter(|p| std::path::Path::new(p).exists())
+                by_toolkit()
             }
         };
         let cccl_inc =
@@ -1349,13 +1456,22 @@ impl CudaDevice {
             )));
         }
 
-        // Fetch PTX.
+        // Fetch the compiled image: native CUBIN (SASS) by default, PTX
+        // under MT_NVRTC_PTX=1. cuModuleLoadData accepts both.
         let ptx = unsafe {
-            let mut ptx_size: usize = 0;
-            nvrtc_check(nvrtcGetPTXSize(prog, &mut ptx_size), "nvrtcGetPTXSize")?;
-            let mut buf = vec![0u8; ptx_size];
-            nvrtc_check(nvrtcGetPTX(prog, buf.as_mut_ptr() as *mut c_char), "nvrtcGetPTX")?;
-            buf
+            if use_ptx {
+                let mut ptx_size: usize = 0;
+                nvrtc_check(nvrtcGetPTXSize(prog, &mut ptx_size), "nvrtcGetPTXSize")?;
+                let mut buf = vec![0u8; ptx_size];
+                nvrtc_check(nvrtcGetPTX(prog, buf.as_mut_ptr() as *mut c_char), "nvrtcGetPTX")?;
+                buf
+            } else {
+                let mut cubin_size: usize = 0;
+                nvrtc_check(nvrtcGetCUBINSize(prog, &mut cubin_size), "nvrtcGetCUBINSize")?;
+                let mut buf = vec![0u8; cubin_size];
+                nvrtc_check(nvrtcGetCUBIN(prog, buf.as_mut_ptr() as *mut c_char), "nvrtcGetCUBIN")?;
+                buf
+            }
         };
         unsafe { nvrtcDestroyProgram(&mut prog) };
 
@@ -1370,8 +1486,7 @@ impl CudaDevice {
 
     /// Allocate `len` bytes of device memory.
     pub fn alloc(&self, len: usize) -> Result<DeviceBuffer<'_>, MetalTileError> {
-        #[allow(unused_assignments)]
-        let mut ptr: CUdeviceptr = 0;
+        self.ensure_current();
         if len == 0 {
             return Ok(DeviceBuffer { ptr: 0, len: 0, _dev: self });
         }
@@ -1380,12 +1495,13 @@ impl CudaDevice {
         // when the pool is disabled, so behavior is unchanged in that mode. nsys
         // showed cuMemAlloc/cuMemFree = 62% of CUDA API time (6110 driver allocs,
         // each ~0.5ms + device-synchronizing → the prefill GPU idle).
-        ptr = self.alloc_raw(len)?;
+        let ptr = self.alloc_raw(len)?;
         Ok(DeviceBuffer { ptr, len, _dev: self })
     }
 
     /// Allocate + upload host bytes in one shot.
     pub fn upload(&self, data: &[u8]) -> Result<DeviceBuffer<'_>, MetalTileError> {
+        self.ensure_current();
         let buf = self.alloc(data.len())?;
         if !data.is_empty() {
             // Enqueue on self.stream (NOT the null stream). Kernels ride self.stream
@@ -1413,6 +1529,7 @@ impl CudaDevice {
 
     /// Copy device memory back into a host buffer.
     pub fn download(&self, buf: &DeviceBuffer, out: &mut [u8]) -> Result<(), MetalTileError> {
+        self.ensure_current();
         let n = out.len().min(buf.len);
         if n == 0 {
             return Ok(());
@@ -1435,6 +1552,7 @@ impl CudaDevice {
     ///
     /// [`free_raw`]: CudaDevice::free_raw
     pub fn alloc_raw(&self, len: usize) -> Result<CUdeviceptr, MetalTileError> {
+        self.ensure_current();
         if len == 0 {
             return Ok(0);
         }
@@ -1448,16 +1566,15 @@ impl CudaDevice {
                 *self.pooled_bytes.lock().unwrap() -= bucket;
                 return Ok(ptr);
             }
-            #[allow(unused_assignments)]
             let mut ptr: CUdeviceptr = 0;
             cu_check(unsafe { cuMemAlloc_v2(&mut ptr, bucket) }, "cuMemAlloc")?;
             return Ok(ptr);
         }
-        // Default (pool off): legacy exact-size pool reuse, else fresh alloc.
+        // Pool off: direct driver alloc. The exact-size pop only ever serves
+        // buffers parked mid-capture by `free_raw_pooled` (see there).
         if let Some(ptr) = self.pool.lock().unwrap().get_mut(&len).and_then(|v| v.pop()) {
             return Ok(ptr);
         }
-        #[allow(unused_assignments)]
         let mut ptr: CUdeviceptr = 0;
         cu_check(unsafe { cuMemAlloc_v2(&mut ptr, len) }, "cuMemAlloc")?;
         Ok(ptr)
@@ -1466,6 +1583,7 @@ impl CudaDevice {
     /// Free a pointer returned by [`alloc_raw`]. No-op on a null pointer.
     /// Eagerly releases to the driver (use [`free_raw_pooled`] on the hot path).
     pub fn free_raw(&self, ptr: CUdeviceptr) {
+        self.ensure_current();
         // Never cuMemFree while a CUDA graph is capturing (the graph captured this
         // pointer → replay would read freed memory). Retain it (leaks for the
         // transient capture; acceptable vs a corrupt graph).
@@ -1478,8 +1596,11 @@ impl CudaDevice {
     /// instead of releasing it to the driver — avoids a synchronous
     /// `cuMemFree`/`cuMemAlloc` round-trip (each a device-wide sync) on the next
     /// same-bucket request. `len` MUST be the value passed to [`alloc_raw`] so
-    /// the bucket re-derives to the one the buffer was allocated as.
+    /// the bucket re-derives to the one the buffer was allocated as. With the
+    /// pool disabled (`METALTILE_POOL_ALLOC_OFF`) this releases to the driver
+    /// directly (except mid-graph-capture, where the pointer must stay live).
     pub fn free_raw_pooled(&self, ptr: CUdeviceptr, len: usize) {
+        self.ensure_current();
         if ptr == 0 {
             return;
         }
@@ -1493,7 +1614,7 @@ impl CudaDevice {
             // buffer pointers, so releasing one to the driver makes replay read
             // freed memory ("unspecified launch failure"). Retain past the cap for
             // the (transient) capture; the pool drains normally afterward.
-            if *parked + bucket > POOL_CAP_BYTES && !self.is_capturing() {
+            if *parked + bucket > pool_cap_bytes() && !self.is_capturing() {
                 drop(parked);
                 unsafe { cuMemFree_v2(ptr) };
                 return;
@@ -1503,9 +1624,15 @@ impl CudaDevice {
             self.pool.lock().unwrap().entry(bucket).or_default().push(ptr);
             return;
         }
-        // Default (pool off): legacy exact-size retain (unbounded; matches the
-        // prior committed behaviour for A/B parity when the flag is unset).
-        self.pool.lock().unwrap().entry(len).or_default().push(ptr);
+        // Pool off: release to the driver — EXCEPT mid-capture (the graph
+        // recorded this pointer; freeing it would make replay read freed
+        // memory). Capture-parked buffers are reused by alloc_raw or freed
+        // on Drop.
+        if self.is_capturing() {
+            self.pool.lock().unwrap().entry(len).or_default().push(ptr);
+            return;
+        }
+        unsafe { cuMemFree_v2(ptr) };
     }
 
     /// Copy host bytes into an existing device allocation (host→device), ASYNC
@@ -1513,6 +1640,7 @@ impl CudaDevice {
     /// without a host-blocking GPU drain. The pinned buffer is reclaimed on the
     /// next `synchronize`/`download` (by then the copy has completed).
     pub fn htod(&self, ptr: CUdeviceptr, bytes: &[u8]) -> Result<(), MetalTileError> {
+        self.ensure_current();
         if bytes.is_empty() {
             return Ok(());
         }
@@ -1536,14 +1664,14 @@ impl CudaDevice {
                 "cuMemcpyHtoDAsync(large)",
             );
         }
-        let pinned =
-            self.pinned_free.lock().unwrap().get_mut(&len).and_then(|v| v.pop()).unwrap_or_else(
-                || {
-                    let mut p: *mut c_void = ptr::null_mut();
-                    unsafe { cuMemAllocHost_v2(&mut p, len) };
-                    p as usize
-                },
-            );
+        let pinned = match self.pinned_free.lock().unwrap().get_mut(&len).and_then(|v| v.pop()) {
+            Some(p) => p,
+            None => {
+                let mut p: *mut c_void = ptr::null_mut();
+                cu_check(unsafe { cuMemAllocHost_v2(&mut p, len) }, "cuMemAllocHost")?;
+                p as usize
+            },
+        };
         unsafe { ptr::copy_nonoverlapping(bytes.as_ptr(), pinned as *mut u8, len) };
         cu_check(
             unsafe { cuMemcpyHtoDAsync_v2(ptr, pinned as *const c_void, len, self.stream) },
@@ -1568,6 +1696,7 @@ impl CudaDevice {
 
     /// Copy device memory back to a host slice (device→host).
     pub fn dtoh(&self, ptr: CUdeviceptr, out: &mut [u8]) -> Result<(), MetalTileError> {
+        self.ensure_current();
         if out.is_empty() {
             return Ok(());
         }
@@ -1591,6 +1720,7 @@ impl CudaDevice {
         block_threads: u32,
         args: &mut [*mut c_void],
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         self.launch(func, [grid_blocks, 1, 1], [block_threads, 1, 1], 0, args)
     }
 
@@ -1605,6 +1735,7 @@ impl CudaDevice {
         shared_bytes: u32,
         args: &mut [*mut c_void],
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         // Opt into the requested dynamic shared size (no-op below the default
         // cap, required above 48KB). Fails *before* launch on archs that cap
         // dynamic smem at 48KB (pre-Volta) instead of a cryptic launch error.
@@ -1643,6 +1774,7 @@ impl CudaDevice {
         shared_bytes: u32,
         args: &mut [*mut c_void],
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         ensure_dynamic_smem(func.func, shared_bytes)?;
         cu_check(
             unsafe {
@@ -1677,6 +1809,7 @@ impl CudaDevice {
         shared_bytes: u32,
         args: &mut [*mut c_void],
     ) -> Result<(), MetalTileError> {
+        self.ensure_current();
         ensure_dynamic_smem(func.func, shared_bytes)?;
         cu_check(
             unsafe {
@@ -1706,6 +1839,7 @@ impl CudaDevice {
     /// megakernel). Caller must run a NO-host-sync (all-device) sequence, then
     /// `end_capture`. THREAD_LOCAL mode scopes capture to this thread's stream.
     pub fn begin_capture(&self) -> Result<(), MetalTileError> {
+        self.ensure_current();
         self.capturing.store(true, std::sync::atomic::Ordering::SeqCst);
         cu_check(
             unsafe { cuStreamBeginCapture_v2(self.stream, CU_STREAM_CAPTURE_MODE_THREAD_LOCAL) },
@@ -1715,22 +1849,31 @@ impl CudaDevice {
 
     /// Finish capture → instantiate an executable graph. Replay with `graph_launch`.
     pub fn end_capture(&self) -> Result<CUgraphExec, MetalTileError> {
+        self.ensure_current();
         let mut graph: CUgraph = ptr::null_mut();
-        cu_check(unsafe { cuStreamEndCapture(self.stream, &mut graph) }, "cuStreamEndCapture")?;
+        let res =
+            cu_check(unsafe { cuStreamEndCapture(self.stream, &mut graph) }, "cuStreamEndCapture");
+        // Clear the flag even on failure — a stuck `capturing` would make
+        // every later free_raw/free_raw_pooled leak "for the capture" forever.
         self.capturing.store(false, std::sync::atomic::Ordering::SeqCst);
+        res?;
         let mut exec: CUgraphExec = ptr::null_mut();
-        cu_check(
+        let inst = cu_check(
             unsafe { cuGraphInstantiateWithFlags(&mut exec, graph, 0) },
             "cuGraphInstantiate",
-        )?;
+        );
         unsafe { cuGraphDestroy(graph) };
+        inst?;
         Ok(exec)
     }
 
     /// Replay a captured decode token: ONE host launch replaces ~390 — no
     /// per-kernel enqueue, no inter-kernel host bubbles. Syncs the stream after.
-    #[allow(clippy::not_unsafe_ptr_arg_deref)] // FFI-handle wrapper, same idiom as the launch fns
+    ///
+    /// `exec` must be a live handle returned by [`Self::end_capture`].
+    #[allow(clippy::not_unsafe_ptr_arg_deref)] // opaque driver handle, not dereferenced host-side
     pub fn graph_launch(&self, exec: CUgraphExec) -> Result<(), MetalTileError> {
+        self.ensure_current();
         cu_check(unsafe { cuGraphLaunch(exec, self.stream) }, "cuGraphLaunch")?;
         self.synchronize()
     }
@@ -1742,14 +1885,55 @@ impl CudaDevice {
     /// `graph_launch` (sync-per-token) incurs, giving the maximum throughput
     /// for the captured graph. Use only for throughput benchmarking — state
     /// (KV cache, SSM state) is overwritten sequentially and not meaningful.
-    #[allow(clippy::not_unsafe_ptr_arg_deref)] // FFI-handle wrapper, same idiom as the launch fns
+    /// `exec` must be a live handle returned by [`Self::end_capture`].
+    #[allow(clippy::not_unsafe_ptr_arg_deref)] // opaque driver handle, not dereferenced host-side
     pub fn graph_launch_batch(&self, exec: CUgraphExec, n: usize) -> Result<(), MetalTileError> {
+        self.ensure_current();
         for _ in 0..n {
             cu_check(unsafe { cuGraphLaunch(exec, self.stream) }, "cuGraphLaunch(batch)")?;
         }
         self.synchronize()
     }
 
+    /// Dispatch-path upload with ZEROED SLACK after the payload. Grid3D
+    /// kernels carry no bounds guard (Metal parity), so a ceil-div launch
+    /// runs tail threads whose loads/stores walk past a buffer's logical
+    /// end. The Metal harness survives that because MTLBuffers are
+    /// page-granular and zero-filled; a recycled CUDA pool bucket instead
+    /// hands those threads garbage (e.g. a gather kernel loading a garbage
+    /// row index → wild address → sticky illegal-memory-access poisoning
+    /// the whole context). Mirror the Metal environment: over-allocate and
+    /// zero-fill so tail threads read benign zeros and their stores land in
+    /// owned slack. Dispatch/test path only — the engine's `alloc_raw`/
+    /// `htod` hot path is unaffected.
+    fn upload_padded(&self, data: &[u8]) -> Result<DeviceBuffer<'_>, MetalTileError> {
+        const DISPATCH_PAD_BYTES: usize = 4096;
+        let buf = self.alloc(data.len() + DISPATCH_PAD_BYTES)?;
+        cu_check(
+            unsafe { cuMemsetD8Async(buf.ptr, 0, buf.len, self.stream) },
+            "cuMemsetD8Async(dispatch pad)",
+        )?;
+        if !data.is_empty() {
+            cu_check(
+                unsafe {
+                    cuMemcpyHtoDAsync_v2(
+                        buf.ptr,
+                        data.as_ptr() as *const c_void,
+                        data.len(),
+                        self.stream,
+                    )
+                },
+                "cuMemcpyHtoDAsync(upload_padded)",
+            )?;
+        }
+        // Sync so the pageable `data` borrow can end (same contract as upload).
+        cu_check(
+            unsafe { cuStreamSynchronize(self.stream) },
+            "cuStreamSynchronize(upload_padded)",
+        )?;
+        Ok(buf)
+    }
+
     /// A compiled+resident kernel: module, function, device buffers, and the
     /// marshalled scalar/pointer args, ready to launch repeatedly without
     /// re-compiling or re-uploading. Produced by [`CudaDevice::prepare`] and
@@ -1775,25 +1959,36 @@ impl CudaDevice {
         }
         let module = self.compile(&src, &format!("{}.cu", kernel.name))?;
         let func = module.function(&kernel.name)?;
-        // Dynamic shared-memory size for this launch geometry.
-        let shared_bytes = cg.shared_bytes(kernel, block[0]) as u32;
-
-        // 2. Allocate + upload each param buffer (in kernel.params order).
-        //    Strided params are not yet supported by the emitter (it errors
-        //    in generate() above), so every param here is Tensor/Scalar.
+        // Dynamic shared-memory size for this launch geometry (per-warp
+        // arrays scale with the TOTAL thread count, x·y·z).
+        let shared_bytes = cg.shared_bytes(kernel, block[0] * block[1] * block[2]) as u32;
+
+        // 2. Marshal each param (in kernel.params order) into arg_slots.
+        //    Tensor/Strided params upload to a device buffer; Scalar params
+        //    stay HOST bytes passed by value through kernelParams — the
+        //    emitted signature takes them by value (`float x`), so a device
+        //    pointer here would be reinterpreted as the scalar's bytes.
         // dev_bufs / dev_ptrs / out_meta are kept index-aligned. Strided
         // params add 2 extra companion buffers, so an output's buffer is NOT
         // at its kernel.params index — out_meta carries the (name,len) for
         // read-back so the alignment is by buffer, not by param.
         let mut dev_bufs: Vec<DeviceBuffer> = Vec::new();
         let mut dev_ptrs: Vec<CUdeviceptr> = Vec::new();
+        let mut scalars: Vec<Vec<u8>> = Vec::new();
+        let mut arg_slots: Vec<ArgSlot> = Vec::new();
         let mut out_meta: Vec<Option<(String, usize)>> = Vec::new();
         for p in &kernel.params {
             let bytes = buffers.get(&p.name).ok_or_else(|| {
                 MetalTileError::Dispatch(format!("missing buffer for param '{}'", p.name))
             })?;
-            let buf = self.upload(bytes)?;
+            if p.kind == metaltile_core::ir::ParamKind::Scalar {
+                scalars.push(bytes.clone());
+                arg_slots.push(ArgSlot::Scalar(scalars.len() - 1));
+                continue;
+            }
+            let buf = self.upload_padded(bytes)?;
             dev_ptrs.push(buf.device_ptr());
+            arg_slots.push(ArgSlot::Buf(dev_ptrs.len() - 1));
             out_meta.push(if p.is_output { Some((p.name.clone(), bytes.len())) } else { None });
             dev_bufs.push(buf);
 
@@ -1807,23 +2002,29 @@ impl CudaDevice {
                         Some(b) => b.clone(),
                         None => synth_strided_meta(&p.shape, suffix == "_strides"),
                     };
-                    let mb = self.upload(&meta)?;
+                    let mb = self.upload_padded(&meta)?;
                     dev_ptrs.push(mb.device_ptr());
+                    arg_slots.push(ArgSlot::Buf(dev_ptrs.len() - 1));
                     out_meta.push(None);
                     dev_bufs.push(mb);
                 }
             }
         }
 
-        // 3. Scalar arg storage: constexprs (signature order) then the
+        // 3. Trailing by-value args: constexprs (signature order) then the
         //    synthetic _n_elems for Elementwise.
-        let mut scalars: Vec<Vec<u8>> = Vec::new();
         for ce in &kernel.constexprs {
             let name = ce.name.name();
-            let bytes = buffers
-                .get(name)
-                .ok_or_else(|| MetalTileError::Dispatch(format!("missing constexpr '{name}'")))?;
-            scalars.push(bytes.clone());
+            // No binding = the constexpr is dead in this variant (e.g. a
+            // shared signature whose pruned body never reads it) — bind a
+            // zero of the declared width, mirroring the Metal dispatch path
+            // which binds a dummy buffer instead of erroring.
+            let bytes = match buffers.get(name) {
+                Some(b) => b.clone(),
+                None => vec![0u8; ce.dtype.size_bytes().max(4)],
+            };
+            scalars.push(bytes);
+            arg_slots.push(ArgSlot::Scalar(scalars.len() - 1));
         }
         if kernel.mode == metaltile_core::ir::KernelMode::Elementwise {
             // Bounds = element count of the first output param.
@@ -1837,9 +2038,19 @@ impl CudaDevice {
                 })
                 .unwrap_or(0);
             scalars.push(n_elems.to_le_bytes().to_vec());
+            arg_slots.push(ArgSlot::Scalar(scalars.len() - 1));
         }
 
-        Ok(Prepared { _module: module, func, dev_bufs, dev_ptrs, scalars, out_meta, shared_bytes })
+        Ok(Prepared {
+            _module: module,
+            func,
+            dev_bufs,
+            dev_ptrs,
+            scalars,
+            arg_slots,
+            out_meta,
+            shared_bytes,
+        })
     }
 
     /// End-to-end generic dispatch: generate CUDA for `kernel`, compile,
@@ -1856,6 +2067,7 @@ impl CudaDevice {
         grid: [u32; 3],
         block: [u32; 3],
     ) -> Result<BTreeMap<String, Vec<u8>>, MetalTileError> {
+        self.ensure_current();
         let prep = self.prepare(kernel, buffers, block)?;
 
         // Launch (with dynamic shared memory).
@@ -1894,6 +2106,7 @@ impl CudaDevice {
         warmup: u32,
         iters: u32,
     ) -> Result<Vec<f64>, MetalTileError> {
+        self.ensure_current();
         let prep = self.prepare(kernel, buffers, block)?;
         let mut args = prep.args();
 
@@ -1913,7 +2126,9 @@ impl CudaDevice {
             let mut samples = Vec::with_capacity(iters as usize);
             for _ in 0..iters {
                 ensure_dynamic_smem(prep.func.func, prep.shared_bytes)?;
-                cu_check(unsafe { cuEventRecord(start, ptr::null_mut()) }, "cuEventRecord(start)")?;
+                // Events + launch all ride self.stream — recording on the null
+                // stream would not order against the non-blocking work stream.
+                cu_check(unsafe { cuEventRecord(start, self.stream) }, "cuEventRecord(start)")?;
                 cu_check(
                     unsafe {
                         cuLaunchKernel(
@@ -1925,14 +2140,14 @@ impl CudaDevice {
                             block[1],
                             block[2],
                             prep.shared_bytes,
-                            ptr::null_mut(),
+                            self.stream,
                             args.as_mut_ptr(),
                             ptr::null_mut(),
                         )
                     },
                     "cuLaunchKernel",
                 )?;
-                cu_check(unsafe { cuEventRecord(stop, ptr::null_mut()) }, "cuEventRecord(stop)")?;
+                cu_check(unsafe { cuEventRecord(stop, self.stream) }, "cuEventRecord(stop)")?;
                 cu_check(unsafe { cuEventSynchronize(stop) }, "cuEventSynchronize")?;
                 let mut ms: f32 = 0.0;
                 cu_check(
@@ -1954,6 +2169,7 @@ impl CudaDevice {
     }
 
     pub fn synchronize(&self) -> Result<(), MetalTileError> {
+        self.ensure_current();
         cu_check(unsafe { cuStreamSynchronize(self.stream) }, "cuStreamSynchronize")?;
         self.reclaim_pinned();
         Ok(())
@@ -1962,6 +2178,42 @@ impl CudaDevice {
 
 impl Drop for CudaDevice {
     fn drop(&mut self) {
+        // Drain the stream before anything is freed — an async H2D copy
+        // still in flight would otherwise DMA from pinned host memory we
+        // are about to cuMemFreeHost.
+        unsafe { cuStreamSynchronize(self.stream) };
+        // Library handles + pinned host memory first: cuCtxDestroy does not
+        // release host pinned allocations, and destroying handles after their
+        // context is gone is undefined.
+        if let Ok(mut h) = self.cublas.lock()
+            && *h != 0
+        {
+            unsafe { cublasDestroy_v2(*h as cublasHandle_t) };
+            *h = 0;
+        }
+        if let Ok(mut lt) = self.cublaslt.lock()
+            && lt.0 != 0
+        {
+            unsafe {
+                cublasLtDestroy(lt.0 as cublasLtHandle_t);
+                if lt.1 != 0 {
+                    cuMemFree_v2(lt.1 as CUdeviceptr);
+                }
+            }
+            *lt = (0, 0);
+        }
+        if let Ok(mut free) = self.pinned_free.lock() {
+            for (_, ptrs) in free.drain() {
+                for p in ptrs {
+                    unsafe { cuMemFreeHost(p as *mut c_void) };
+                }
+            }
+        }
+        if let Ok(mut inflight) = self.pinned_inflight.lock() {
+            for (p, _) in inflight.drain(..) {
+                unsafe { cuMemFreeHost(p as *mut c_void) };
+            }
+        }
         if let Ok(mut pool) = self.pool.lock() {
             for (_, ptrs) in pool.drain() {
                 for p in ptrs {
@@ -1972,8 +2224,14 @@ impl Drop for CudaDevice {
         if let Ok(mut parked) = self.pooled_bytes.lock() {
             *parked = 0;
         }
+        if !self.stream.is_null() {
+            unsafe { cuStreamDestroy_v2(self.stream) };
+        }
         if !self.ctx.is_null() {
-            unsafe { cuCtxDestroy_v2(self.ctx) };
+            // Primary context: release the refcount taken in create() —
+            // destroying it would tear the context down under any other
+            // holder in the process.
+            unsafe { cuDevicePrimaryCtxRelease_v2(self.dev) };
         }
     }
 }
diff --git a/crates/metaltile-runtime/src/device/cuda/nvfp4_moe.rs b/crates/metaltile-runtime/src/device/cuda/nvfp4_moe.rs
new file mode 100644
index 00000000..5721dd51
--- /dev/null
+++ b/crates/metaltile-runtime/src/device/cuda/nvfp4_moe.rs
@@ -0,0 +1,1130 @@
+//! Copyright 2026 0xClandestine, Ekryski, TheTom, Ambisphaeric
+//! SPDX-License-Identifier: Apache-2.0
+//! NVFP4 / FP8 cutlass grouped-MoE GEMM entry points for `CudaDevice`.
+//!
+//! Carved out of the kitchen-sink CUDA work and re-architected onto the current
+//! backend: additive `impl CudaDevice` methods wrapping the AOT-compiled cutlass
+//! grouped kernels (`cuda/cutlass_moe_fp4.cu`, gated behind `CUTLASS_DIR` /
+//! `cfg(have_cutlass)` in build.rs) plus the cuBLASLt fp4/fp8 GEMM escape hatch.
+//! Kept in a dedicated module so the NVFP4 surface stays separate from the core
+//! device lifecycle in `mod.rs`.
+//!
+//! Two lints are relaxed module-wide to match the C side: `too_many_arguments`
+//! (the GEMM/MoE wrappers mirror the cutlass / cuBLASLt C ABI parameter lists)
+//! and `non_snake_case` (a few entry points embed the cutlass kernel variant
+//! tags FUSEDACT / AMAX verbatim so the Rust name maps 1:1 to the CUDA variant).
+#![allow(clippy::too_many_arguments, non_snake_case)]
+
+use std::{
+    os::raw::{c_int, c_void},
+    ptr,
+};
+
+use super::{CUBLASLT_WORKSPACE_BYTES, CUdeviceptr, CudaDevice, cu_check, ffi::*};
+use crate::error::MetalTileError;
+
+impl CudaDevice {
+    /// Lazily create the cublasLt handle + fixed device workspace. Returns
+    /// `(handle, workspace_ptr)`.
+    fn lt_beta_zero_ptr(&self) -> Result<CUdeviceptr, MetalTileError> {
+        let mut g = self.lt_beta_zero.lock().unwrap();
+        if *g == 0 {
+            let p = self.alloc_raw(4)?;
+            cu_check(
+                unsafe { cuMemsetD8Async(p, 0, 4, self.stream) },
+                "cuMemsetD8Async(lt_beta_zero)",
+            )?;
+            *g = p as usize;
+        }
+        Ok(*g as CUdeviceptr)
+    }
+
+    /// CUTLASS grouped block-scaled NVFP4 MoE GEMM (AOT-linked, sm_120a/121a).
+    /// `a` = packed e2m1 sorted-token activations `[mt, K/2]` bytes, `sfa` =
+    /// per-group ue4m3 scale pool (group g's blob at byte `sfa_off[g]`, laid out
+    /// for the group-LOCAL row), `w`/`sfw` = packed e2m1 + scale expert slabs
+    /// (`[n_exp, N*K/2]` / `[n_exp, ceil(N/128)*512*ceil(K/64)]` bytes), `c` =
+    /// f16 out `[mt, N]`. `group_rows`/`expert_ids`/`sfa_off` are HOST slices.
+    /// `out[t,n] = Σ_k a[t,k]·w[eid][n,k]` with per-16-block scales on both
+    /// operands. Errors if the runtime was built without CUTLASS.
+    #[allow(clippy::too_many_arguments)] // mirrors the C entry point's signature
+    pub fn moe_grouped_cutlass_fp4(
+        &self,
+        a: CUdeviceptr,
+        sfa: CUdeviceptr,
+        w: CUdeviceptr,
+        sfw: CUdeviceptr,
+        c: CUdeviceptr,
+        group_rows: &[i32],
+        expert_ids: &[i32],
+        sfa_off: &[i64],
+        alpha_vec: CUdeviceptr, // device f32[n_groups] per-group scales, 0 = none
+        n: usize,
+        k: usize,
+    ) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn moe_grouped_gemm_cutlass_fp4(
+                    a: *const c_void,
+                    sfa: *const c_void,
+                    w: *const c_void,
+                    sfw: *const c_void,
+                    c: *mut c_void,
+                    group_rows: *const c_int,
+                    expert_ids: *const c_int,
+                    sfa_off: *const i64,
+                    alpha_vec: *const c_void,
+                    n_groups: c_int,
+                    n: c_int,
+                    k: c_int,
+                    stream: *mut c_void,
+                ) -> c_int;
+            }
+            if group_rows.len() != expert_ids.len() || group_rows.len() != sfa_off.len() {
+                return Err(MetalTileError::Dispatch(
+                    "moe_grouped_cutlass_fp4: group_rows/expert_ids/sfa_off len mismatch".into(),
+                ));
+            }
+            let r = unsafe {
+                moe_grouped_gemm_cutlass_fp4(
+                    a as *const c_void,
+                    sfa as *const c_void,
+                    w as *const c_void,
+                    sfw as *const c_void,
+                    c as *mut c_void,
+                    group_rows.as_ptr(),
+                    expert_ids.as_ptr(),
+                    sfa_off.as_ptr(),
+                    alpha_vec as *const c_void,
+                    group_rows.len() as c_int,
+                    n as c_int,
+                    k as c_int,
+                    self.stream as *mut c_void,
+                )
+            };
+            if r != 0 {
+                return Err(MetalTileError::Dispatch(format!(
+                    "moe_grouped_gemm_cutlass_fp4 failed: code {r}"
+                )));
+            }
+            Ok(())
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (a, sfa, w, sfw, c, group_rows, expert_ids, sfa_off, alpha_vec, n, k);
+            Err(MetalTileError::Dispatch(
+                "moe_grouped_cutlass_fp4: runtime built without CUTLASS (set CUTLASS_DIR)".into(),
+            ))
+        }
+    }
+
+    /// W4A8 grouped GEMM (fp8 e4m3 acts × fp4 e2m1 weights, mxf8f6f4). Mirrors
+    /// moe_grouped_cutlass_fp4 + sfb_exp_bytes (per-expert SFB stride from w4a8_packw_run).
+    pub fn moe_grouped_cutlass_w4a8(
+        &self,
+        a: CUdeviceptr,
+        sfa: CUdeviceptr,
+        w: CUdeviceptr,
+        sfw: CUdeviceptr,
+        c: CUdeviceptr,
+        group_rows: &[i32],
+        expert_ids: &[i32],
+        sfa_off: &[i64],
+        sfb_exp_bytes: i64,
+        alpha_vec: CUdeviceptr,
+        n: usize,
+        k: usize,
+    ) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn moe_grouped_gemm_w4a8(
+                    a: *const c_void,
+                    sfa: *const c_void,
+                    w: *const c_void,
+                    sfw: *const c_void,
+                    c: *mut c_void,
+                    group_rows: *const c_int,
+                    expert_ids: *const c_int,
+                    sfa_off: *const i64,
+                    sfb_exp_bytes: i64,
+                    alpha_vec: *const c_void,
+                    n_groups: c_int,
+                    n: c_int,
+                    k: c_int,
+                    stream: *mut c_void,
+                ) -> c_int;
+            }
+            if group_rows.len() != expert_ids.len() || group_rows.len() != sfa_off.len() {
+                return Err(MetalTileError::Dispatch(
+                    "moe_grouped_cutlass_w4a8: len mismatch".into(),
+                ));
+            }
+            let r = unsafe {
+                moe_grouped_gemm_w4a8(
+                    a as *const c_void,
+                    sfa as *const c_void,
+                    w as *const c_void,
+                    sfw as *const c_void,
+                    c as *mut c_void,
+                    group_rows.as_ptr(),
+                    expert_ids.as_ptr(),
+                    sfa_off.as_ptr(),
+                    sfb_exp_bytes,
+                    alpha_vec as *const c_void,
+                    group_rows.len() as c_int,
+                    n as c_int,
+                    k as c_int,
+                    self.stream as *mut c_void,
+                )
+            };
+            if r != 0 {
+                return Err(MetalTileError::Dispatch(format!("moe_grouped_gemm_w4a8 failed: {r}")));
+            }
+            Ok(())
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (
+                a,
+                sfa,
+                w,
+                sfw,
+                c,
+                group_rows,
+                expert_ids,
+                sfa_off,
+                sfb_exp_bytes,
+                alpha_vec,
+                n,
+                k,
+            );
+            Err(MetalTileError::Dispatch("moe_grouped_cutlass_w4a8: built without CUTLASS".into()))
+        }
+    }
+
+    /// W4A8 fp8 act-quant: x[mt,K] half -> out[mt,K] e4m3 + sf ue8m0 (per-32, group-aware SFA).
+    /// group_rows / sfa_off are HOST arrays (per-group SFA section offsets, atom-padded).
+    pub fn w4a8_actquant_run(
+        &self,
+        x: CUdeviceptr,
+        out: CUdeviceptr,
+        sf: CUdeviceptr,
+        group_rows: &[i32],
+        sfa_off: &[i64],
+        n: usize,
+        k: usize,
+    ) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn w4a8_actquant(
+                    x: *const c_void,
+                    out: *mut c_void,
+                    sf: *mut c_void,
+                    group_rows: *const c_int,
+                    sfa_off: *const ::std::os::raw::c_longlong,
+                    n_groups: c_int,
+                    n: c_int,
+                    k: c_int,
+                    stream: *mut c_void,
+                );
+            }
+            let n_groups = group_rows.len() as c_int;
+            unsafe {
+                w4a8_actquant(
+                    x as *const c_void,
+                    out as *mut c_void,
+                    sf as *mut c_void,
+                    group_rows.as_ptr(),
+                    sfa_off.as_ptr() as *const ::std::os::raw::c_longlong,
+                    n_groups,
+                    n as c_int,
+                    k as c_int,
+                    self.stream as *mut c_void,
+                );
+            }
+            Ok(())
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (x, out, sf, group_rows, sfa_off, n, k);
+            Err(MetalTileError::Dispatch("w4a8_actquant: built without CUTLASS".into()))
+        }
+    }
+
+    /// W4A8 mxfp4 weight pack: w[n_exp,N,K] half -> outp e2m1 + sf ue8m0 (per-32, SFB).
+    /// Returns the per-expert SFB element count (bytes; size the SF buffer + stride).
+    pub fn w4a8_packw_run(
+        &self,
+        w: CUdeviceptr,
+        outp: CUdeviceptr,
+        sf: CUdeviceptr,
+        n_exp: usize,
+        n: usize,
+        k: usize,
+    ) -> Result<i64, MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn w4a8_packw(
+                    w: *const c_void,
+                    outp: *mut c_void,
+                    sf: *mut c_void,
+                    n_exp: c_int,
+                    n: c_int,
+                    k: c_int,
+                    stream: *mut c_void,
+                ) -> i64;
+            }
+            Ok(unsafe {
+                w4a8_packw(
+                    w as *const c_void,
+                    outp as *mut c_void,
+                    sf as *mut c_void,
+                    n_exp as c_int,
+                    n as c_int,
+                    k as c_int,
+                    self.stream as *mut c_void,
+                )
+            })
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (w, outp, sf, n_exp, n, k);
+            Err(MetalTileError::Dispatch("w4a8_packw: built without CUTLASS".into()))
+        }
+    }
+
+    pub fn moe_grouped_cutlass_w8a8(
+        &self,
+        a: CUdeviceptr,
+        sfa: CUdeviceptr,
+        w: CUdeviceptr,
+        sfw: CUdeviceptr,
+        c: CUdeviceptr,
+        group_rows: &[i32],
+        expert_ids: &[i32],
+        sfa_off: &[i64],
+        sfb_exp_bytes: i64,
+        alpha_vec: CUdeviceptr,
+        n: usize,
+        k: usize,
+    ) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn moe_grouped_gemm_w8a8(
+                    a: *const c_void,
+                    sfa: *const c_void,
+                    w: *const c_void,
+                    sfw: *const c_void,
+                    c: *mut c_void,
+                    group_rows: *const c_int,
+                    expert_ids: *const c_int,
+                    sfa_off: *const i64,
+                    sfb_exp_bytes: i64,
+                    alpha_vec: *const c_void,
+                    n_groups: c_int,
+                    n: c_int,
+                    k: c_int,
+                    stream: *mut c_void,
+                ) -> c_int;
+            }
+            if group_rows.len() != expert_ids.len() || group_rows.len() != sfa_off.len() {
+                return Err(MetalTileError::Dispatch(
+                    "moe_grouped_cutlass_w8a8: len mismatch".into(),
+                ));
+            }
+            let r = unsafe {
+                moe_grouped_gemm_w8a8(
+                    a as *const c_void,
+                    sfa as *const c_void,
+                    w as *const c_void,
+                    sfw as *const c_void,
+                    c as *mut c_void,
+                    group_rows.as_ptr(),
+                    expert_ids.as_ptr(),
+                    sfa_off.as_ptr(),
+                    sfb_exp_bytes,
+                    alpha_vec as *const c_void,
+                    group_rows.len() as c_int,
+                    n as c_int,
+                    k as c_int,
+                    self.stream as *mut c_void,
+                )
+            };
+            if r != 0 {
+                return Err(MetalTileError::Dispatch(format!("moe_grouped_gemm_w8a8 failed: {r}")));
+            }
+            Ok(())
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (
+                a,
+                sfa,
+                w,
+                sfw,
+                c,
+                group_rows,
+                expert_ids,
+                sfa_off,
+                sfb_exp_bytes,
+                alpha_vec,
+                n,
+                k,
+            );
+            Err(MetalTileError::Dispatch("moe_grouped_cutlass_w8a8: built without CUTLASS".into()))
+        }
+    }
+
+    /// W4A8 fp8 act-quant: x[mt,K] half -> out[mt,K] e4m3 + sf ue8m0 (per-32, group-aware SFA).
+    /// group_rows / sfa_off are HOST arrays (per-group SFA section offsets, atom-padded).
+    pub fn w8a8_actquant_run(
+        &self,
+        x: CUdeviceptr,
+        out: CUdeviceptr,
+        sf: CUdeviceptr,
+        group_rows: &[i32],
+        sfa_off: &[i64],
+        n: usize,
+        k: usize,
+    ) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn w8a8_actquant(
+                    x: *const c_void,
+                    out: *mut c_void,
+                    sf: *mut c_void,
+                    group_rows: *const c_int,
+                    sfa_off: *const ::std::os::raw::c_longlong,
+                    n_groups: c_int,
+                    n: c_int,
+                    k: c_int,
+                    stream: *mut c_void,
+                );
+            }
+            let n_groups = group_rows.len() as c_int;
+            unsafe {
+                w8a8_actquant(
+                    x as *const c_void,
+                    out as *mut c_void,
+                    sf as *mut c_void,
+                    group_rows.as_ptr(),
+                    sfa_off.as_ptr() as *const ::std::os::raw::c_longlong,
+                    n_groups,
+                    n as c_int,
+                    k as c_int,
+                    self.stream as *mut c_void,
+                );
+            }
+            Ok(())
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (x, out, sf, group_rows, sfa_off, n, k);
+            Err(MetalTileError::Dispatch("w8a8_actquant: built without CUTLASS".into()))
+        }
+    }
+
+    /// W4A8 mxfp4 weight pack: w[n_exp,N,K] half -> outp e2m1 + sf ue8m0 (per-32, SFB).
+    /// Returns the per-expert SFB element count (bytes; size the SF buffer + stride).
+    pub fn w8a8_packw_run(
+        &self,
+        w: CUdeviceptr,
+        outp: CUdeviceptr,
+        sf: CUdeviceptr,
+        n_exp: usize,
+        n: usize,
+        k: usize,
+    ) -> Result<i64, MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn w8a8_packw(
+                    w: *const c_void,
+                    outp: *mut c_void,
+                    sf: *mut c_void,
+                    n_exp: c_int,
+                    n: c_int,
+                    k: c_int,
+                    stream: *mut c_void,
+                ) -> i64;
+            }
+            Ok(unsafe {
+                w8a8_packw(
+                    w as *const c_void,
+                    outp as *mut c_void,
+                    sf as *mut c_void,
+                    n_exp as c_int,
+                    n as c_int,
+                    k as c_int,
+                    self.stream as *mut c_void,
+                )
+            })
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (w, outp, sf, n_exp, n, k);
+            Err(MetalTileError::Dispatch("w8a8_packw: built without CUTLASS".into()))
+        }
+    }
+
+    /// Persistent-handle variant of the CUTLASS grouped NVFP4 GEMM: descriptors
+    /// are derived ON DEVICE from the group-offsets buffer each call (one tiny
+    /// fill kernel + gemm.run) — no host build, no per-call allocs, graph-safe.
+    /// `prepare` once per (weight slab, n_groups, N, K); `run` per call.
+    #[allow(clippy::too_many_arguments)]
+    pub fn moe_grouped_cutlass_fp4_prepare(
+        &self,
+        w: CUdeviceptr,
+        sfw: CUdeviceptr,
+        alpha_vec: CUdeviceptr,
+        n_groups: usize,
+        n: usize,
+        k: usize,
+        max_m_total: usize,
+    ) -> Result<u64, MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn moe_grouped_gemm_cutlass_fp4_prepare(
+                    b: *const c_void,
+                    sfb: *const c_void,
+                    alpha_vec: *const c_void,
+                    n_groups: c_int,
+                    n: c_int,
+                    k: c_int,
+                    max_m_total: c_int,
+                ) -> *mut c_void;
+            }
+            let h = unsafe {
+                moe_grouped_gemm_cutlass_fp4_prepare(
+                    w as *const c_void,
+                    sfw as *const c_void,
+                    alpha_vec as *const c_void,
+                    n_groups as c_int,
+                    n as c_int,
+                    k as c_int,
+                    max_m_total as c_int,
+                )
+            };
+            if h.is_null() {
+                return Err(MetalTileError::Dispatch("cutlass_fp4_prepare failed".into()));
+            }
+            Ok(h as u64)
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (w, sfw, alpha_vec, n_groups, n, k, max_m_total);
+            Err(MetalTileError::Dispatch("built without CUTLASS (set CUTLASS_DIR)".into()))
+        }
+    }
+
+    /// Per-call run against a prepared handle. `off_dev` = device u32
+    /// `[n_groups+1]` row offsets (the router's expert offsets).
+    pub fn moe_grouped_cutlass_fp4_run(
+        &self,
+        handle: u64,
+        a: CUdeviceptr,
+        sfa: CUdeviceptr,
+        d_out: CUdeviceptr,
+        off_dev: CUdeviceptr,
+    ) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn moe_grouped_gemm_cutlass_fp4_run(
+                    handle: *mut c_void,
+                    a: *const c_void,
+                    sfa: *const c_void,
+                    d: *mut c_void,
+                    off_dev: *const c_void,
+                    stream: *mut c_void,
+                ) -> c_int;
+            }
+            let r = unsafe {
+                moe_grouped_gemm_cutlass_fp4_run(
+                    handle as *mut c_void,
+                    a as *const c_void,
+                    sfa as *const c_void,
+                    d_out as *mut c_void,
+                    off_dev as *const c_void,
+                    self.stream as *mut c_void,
+                )
+            };
+            if r != 0 {
+                return Err(MetalTileError::Dispatch(format!("cutlass_fp4_run failed: code {r}")));
+            }
+            Ok(())
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (handle, a, sfa, d_out, off_dev);
+            Err(MetalTileError::Dispatch("built without CUTLASS (set CUTLASS_DIR)".into()))
+        }
+    }
+
+    /// Fused-activation variant of [`moe_grouped_cutlass_fp4_prepare`]: GEMM1's
+    /// epilogue folds in relu²·(1/256) + NVFP4 block-scale-quant, so the GEMM
+    /// emits e2m1 D + ue4m3 SFD directly (the down-GEMM's input). `norm_constant`
+    /// = device f32 ptr to the static gs (1/256). Gated behind NEMOTRON_FUSE_UPACT.
+    #[allow(clippy::too_many_arguments)]
+    pub fn moe_grouped_cutlass_fp4_FUSEDACT_prepare(
+        &self,
+        w: CUdeviceptr,
+        sfw: CUdeviceptr,
+        alpha_vec: CUdeviceptr,
+        norm_constant: CUdeviceptr,
+        n_groups: usize,
+        n: usize,
+        k: usize,
+        max_m_total: usize,
+    ) -> Result<u64, MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn moe_grouped_gemm_cutlass_fp4_FUSEDACT_prepare(
+                    b: *const c_void,
+                    sfb: *const c_void,
+                    alpha_vec: *const c_void,
+                    norm_constant: *const c_void,
+                    n_groups: c_int,
+                    n: c_int,
+                    k: c_int,
+                    max_m_total: c_int,
+                ) -> *mut c_void;
+            }
+            let h = unsafe {
+                moe_grouped_gemm_cutlass_fp4_FUSEDACT_prepare(
+                    w as *const c_void,
+                    sfw as *const c_void,
+                    alpha_vec as *const c_void,
+                    norm_constant as *const c_void,
+                    n_groups as c_int,
+                    n as c_int,
+                    k as c_int,
+                    max_m_total as c_int,
+                )
+            };
+            if h.is_null() {
+                return Err(MetalTileError::Dispatch("cutlass_fp4_FUSEDACT_prepare failed".into()));
+            }
+            Ok(h as u64)
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (w, sfw, alpha_vec, norm_constant, n_groups, n, k, max_m_total);
+            Err(MetalTileError::Dispatch("built without CUTLASS (set CUTLASS_DIR)".into()))
+        }
+    }
+
+    /// Per-call run against a fused-activation handle. Emits `d_out` (e2m1 packed
+    /// `[mt, N/2]`) + `sfd_out` (ue4m3 swizzled SF blob, same layout the plain
+    /// down-GEMM reads as SFA). `off_dev` = device u32 `[n_groups+1]` row offsets.
+    #[allow(clippy::too_many_arguments)]
+    pub fn moe_grouped_cutlass_fp4_FUSEDACT_run(
+        &self,
+        handle: u64,
+        a: CUdeviceptr,
+        sfa: CUdeviceptr,
+        d_out: CUdeviceptr,
+        sfd_out: CUdeviceptr,
+        off_dev: CUdeviceptr,
+    ) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn moe_grouped_gemm_cutlass_fp4_FUSEDACT_run(
+                    handle: *mut c_void,
+                    a: *const c_void,
+                    sfa: *const c_void,
+                    d: *mut c_void,
+                    sfd: *mut c_void,
+                    off_dev: *const c_void,
+                    stream: *mut c_void,
+                ) -> c_int;
+            }
+            let r = unsafe {
+                moe_grouped_gemm_cutlass_fp4_FUSEDACT_run(
+                    handle as *mut c_void,
+                    a as *const c_void,
+                    sfa as *const c_void,
+                    d_out as *mut c_void,
+                    sfd_out as *mut c_void,
+                    off_dev as *const c_void,
+                    self.stream as *mut c_void,
+                )
+            };
+            if r != 0 {
+                return Err(MetalTileError::Dispatch(format!(
+                    "cutlass_fp4_FUSEDACT_run failed: code {r}"
+                )));
+            }
+            Ok(())
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (handle, a, sfa, d_out, sfd_out, off_dev);
+            Err(MetalTileError::Dispatch("built without CUTLASS (set CUTLASS_DIR)".into()))
+        }
+    }
+
+    /// Amax-in-epilogue variant of [`moe_grouped_cutlass_fp4_prepare`]: GEMM1
+    /// emits f16 D = relu²·(1/256) (the a2 / up_out) AND a per-group amax of
+    /// those activated values into `d_amax` (device `f32[n_groups]`, persistent
+    /// → graph-safe; max over groups = the per-tensor global for the down-quant).
+    /// The down-quant then reads a2 ONCE (no separate amax scan). Gated behind
+    /// NEMOTRON_AMAX_EPI. Output is bit-identical to the 2-pass relu2+amax+quant.
+    #[allow(clippy::too_many_arguments)]
+    pub fn moe_grouped_cutlass_fp4_AMAX_prepare(
+        &self,
+        w: CUdeviceptr,
+        sfw: CUdeviceptr,
+        d_amax: CUdeviceptr,
+        n_groups: usize,
+        n: usize,
+        k: usize,
+        max_m_total: usize,
+    ) -> Result<u64, MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn moe_grouped_gemm_cutlass_fp4_AMAX_prepare(
+                    b: *const c_void,
+                    sfb: *const c_void,
+                    d_amax: *const c_void,
+                    n_groups: c_int,
+                    n: c_int,
+                    k: c_int,
+                    max_m_total: c_int,
+                ) -> *mut c_void;
+            }
+            let h = unsafe {
+                moe_grouped_gemm_cutlass_fp4_AMAX_prepare(
+                    w as *const c_void,
+                    sfw as *const c_void,
+                    d_amax as *const c_void,
+                    n_groups as c_int,
+                    n as c_int,
+                    k as c_int,
+                    max_m_total as c_int,
+                )
+            };
+            if h.is_null() {
+                return Err(MetalTileError::Dispatch("cutlass_fp4_AMAX_prepare failed".into()));
+            }
+            Ok(h as u64)
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (w, sfw, d_amax, n_groups, n, k, max_m_total);
+            Err(MetalTileError::Dispatch("built without CUTLASS (set CUTLASS_DIR)".into()))
+        }
+    }
+
+    /// Per-call run against an amax-in-epilogue handle. Emits `d_out` (f16 a2,
+    /// `[mt, N]`) + writes the per-group amax into the handle's `d_amax`. The
+    /// caller MUST zero `d_amax` before each call (the build uses
+    /// `-DCUTLASS_SKIP_REDUCTION_INIT=1`, so the kernel does not self-init it).
+    /// `off_dev` = device u32 `[n_groups+1]` row offsets.
+    pub fn moe_grouped_cutlass_fp4_AMAX_run(
+        &self,
+        handle: u64,
+        a: CUdeviceptr,
+        sfa: CUdeviceptr,
+        d_out: CUdeviceptr,
+        off_dev: CUdeviceptr,
+    ) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        #[cfg(have_cutlass)]
+        {
+            unsafe extern "C" {
+                fn moe_grouped_gemm_cutlass_fp4_AMAX_run(
+                    handle: *mut c_void,
+                    a: *const c_void,
+                    sfa: *const c_void,
+                    d: *mut c_void,
+                    off_dev: *const c_void,
+                    stream: *mut c_void,
+                ) -> c_int;
+            }
+            let r = unsafe {
+                moe_grouped_gemm_cutlass_fp4_AMAX_run(
+                    handle as *mut c_void,
+                    a as *const c_void,
+                    sfa as *const c_void,
+                    d_out as *mut c_void,
+                    off_dev as *const c_void,
+                    self.stream as *mut c_void,
+                )
+            };
+            if r != 0 {
+                return Err(MetalTileError::Dispatch(format!(
+                    "cutlass_fp4_AMAX_run failed: code {r}"
+                )));
+            }
+            Ok(())
+        }
+        #[cfg(not(have_cutlass))]
+        {
+            let _ = (handle, a, sfa, d_out, off_dev);
+            Err(MetalTileError::Dispatch("built without CUTLASS (set CUTLASS_DIR)".into()))
+        }
+    }
+
+    /// Block-scaled NVFP4 GEMM via cuBLASLt (Blackwell tensor cores):
+    ///   `out[m,n](f16) = X[m,k](e2m1, vec16 ue4m3 scales) · W[n,k]^T(e2m1, vec16 ue4m3)`
+    /// Operands are packed 2 elems/byte row-major; scale tensors use the
+    /// 512-byte-block swizzled layout (one ue4m3 per 16 elements along K):
+    ///   sf_off(r, kb) = (r/128)*512*ceil(KB/4) + (kb/4)*512 + (r%32)*16
+    ///                 + ((r/32)%4)*4 + (kb%4)
+    /// Measured 240-306 TFLOP/s on GB10 at dense-projection shapes — ~4x the
+    /// f16 path — with max_rel 5e-4 vs an exact dequant reference.
+    #[allow(clippy::too_many_arguments)]
+    pub fn gemm_cublaslt_fp4(
+        &self,
+        x: CUdeviceptr,    // [m, k/2] packed e2m1 activation
+        x_sf: CUdeviceptr, // activation scales, swizzled
+        w: CUdeviceptr,    // [n, k/2] packed e2m1 weight
+        w_sf: CUdeviceptr, // weight scales, swizzled
+        out: CUdeviceptr,  // [m, n] row-major result (f16, or f32 when out_f32)
+        m: usize,
+        n: usize,
+        k: usize,
+        out_f32: bool,
+        d_scale: CUdeviceptr, // device f32 applied to D (per-tensor global fold); 0 = none
+    ) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        let (lt, workspace) = self.cublaslt_ctx()?;
+        let alpha: f32 = 1.0;
+        let beta: f32 = 0.0;
+        let use_dev_alpha = d_scale != 0;
+        let beta_zero = if use_dev_alpha { self.lt_beta_zero_ptr()? } else { 0 };
+        unsafe {
+            let mut desc: cublasLtMatmulDesc_t = ptr::null_mut();
+            let s = cublasLtMatmulDescCreate(&mut desc, CUBLAS_COMPUTE_32F, CUDA_R_32F);
+            if s != CUBLAS_STATUS_SUCCESS {
+                return Err(MetalTileError::Dispatch(format!("cublasLtMatmulDescCreate: {s}")));
+            }
+            if use_dev_alpha {
+                let pm: c_int = CUBLASLT_POINTER_MODE_DEVICE;
+                cublasLtMatmulDescSetAttribute(
+                    desc,
+                    CUBLASLT_MATMUL_DESC_POINTER_MODE,
+                    &pm as *const c_int as *const c_void,
+                    std::mem::size_of::<c_int>(),
+                );
+            }
+            let opt = CUBLAS_OP_T;
+            let opn = CUBLAS_OP_N;
+            cublasLtMatmulDescSetAttribute(
+                desc,
+                CUBLASLT_MATMUL_DESC_TRANSA,
+                &opt as *const c_int as *const c_void,
+                std::mem::size_of::<c_int>(),
+            );
+            cublasLtMatmulDescSetAttribute(
+                desc,
+                CUBLASLT_MATMUL_DESC_TRANSB,
+                &opn as *const c_int as *const c_void,
+                std::mem::size_of::<c_int>(),
+            );
+            // Block-scale modes + pointers. With the col-major reorder A=W and
+            // B=X, so the A scale is the WEIGHT scale and B the activation.
+            let sm: c_int = CUBLASLT_MATMUL_MATRIX_SCALE_VEC16_UE4M3;
+            cublasLtMatmulDescSetAttribute(
+                desc,
+                CUBLASLT_MATMUL_DESC_A_SCALE_MODE,
+                &sm as *const c_int as *const c_void,
+                std::mem::size_of::<c_int>(),
+            );
+            cublasLtMatmulDescSetAttribute(
+                desc,
+                CUBLASLT_MATMUL_DESC_B_SCALE_MODE,
+                &sm as *const c_int as *const c_void,
+                std::mem::size_of::<c_int>(),
+            );
+            cublasLtMatmulDescSetAttribute(
+                desc,
+                CUBLASLT_MATMUL_DESC_A_SCALE_POINTER,
+                &w_sf as *const CUdeviceptr as *const c_void,
+                std::mem::size_of::<CUdeviceptr>(),
+            );
+            cublasLtMatmulDescSetAttribute(
+                desc,
+                CUBLASLT_MATMUL_DESC_B_SCALE_POINTER,
+                &x_sf as *const CUdeviceptr as *const c_void,
+                std::mem::size_of::<CUdeviceptr>(),
+            );
+            // A=W [n,k] rm == [k,n] cm (ld=k); B=X [m,k] rm == [k,m] cm (ld=k);
+            // D [m,n] rm == [n,m] cm (ld=n). lds are in ELEMENTS (4-bit ok: k%32==0).
+            let mut a_l: cublasLtMatrixLayout_t = ptr::null_mut();
+            let mut b_l: cublasLtMatrixLayout_t = ptr::null_mut();
+            let mut d_l: cublasLtMatrixLayout_t = ptr::null_mut();
+            cublasLtMatrixLayoutCreate(&mut a_l, CUDA_R_4F_E2M1, k as u64, n as u64, k as i64);
+            cublasLtMatrixLayoutCreate(&mut b_l, CUDA_R_4F_E2M1, k as u64, m as u64, k as i64);
+            cublasLtMatrixLayoutCreate(
+                &mut d_l,
+                if out_f32 { CUDA_R_32F } else { CUDA_R_16F },
+                n as u64,
+                m as u64,
+                n as i64,
+            );
+
+            let mut pref: cublasLtMatmulPreference_t = ptr::null_mut();
+            cublasLtMatmulPreferenceCreate(&mut pref);
+            let ws_bytes: usize = CUBLASLT_WORKSPACE_BYTES;
+            cublasLtMatmulPreferenceSetAttribute(
+                pref,
+                CUBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES,
+                &ws_bytes as *const usize as *const c_void,
+                std::mem::size_of::<usize>(),
+            );
+            let red_mask: u32 = CUBLASLT_REDUCTION_SCHEME_NONE;
+            cublasLtMatmulPreferenceSetAttribute(
+                pref,
+                CUBLASLT_MATMUL_PREF_REDUCTION_SCHEME_MASK,
+                &red_mask as *const u32 as *const c_void,
+                std::mem::size_of::<u32>(),
+            );
+
+            let mut result = cublasLtMatmulHeuristicResult_t::default();
+            let mut returned: c_int = 0;
+            let hs = cublasLtMatmulAlgoGetHeuristic(
+                lt,
+                desc,
+                a_l,
+                b_l,
+                d_l,
+                d_l,
+                pref,
+                1,
+                &mut result,
+                &mut returned,
+            );
+            if hs != CUBLAS_STATUS_SUCCESS || returned < 1 {
+                cublasLtMatmulPreferenceDestroy(pref);
+                cublasLtMatrixLayoutDestroy(a_l);
+                cublasLtMatrixLayoutDestroy(b_l);
+                cublasLtMatrixLayoutDestroy(d_l);
+                cublasLtMatmulDescDestroy(desc);
+                return Err(MetalTileError::Dispatch(format!(
+                    "cublasLt-fp4: no algo (m={m} n={n} k={k} status={hs} returned={returned})"
+                )));
+            }
+            let alpha_arg: *const c_void = if use_dev_alpha {
+                d_scale as usize as *const c_void
+            } else {
+                &alpha as *const f32 as *const c_void
+            };
+            let beta_arg: *const c_void = if use_dev_alpha {
+                beta_zero as usize as *const c_void
+            } else {
+                &beta as *const f32 as *const c_void
+            };
+            let mm = cublasLtMatmul(
+                lt,
+                desc,
+                alpha_arg,
+                w,
+                a_l,
+                x,
+                b_l,
+                beta_arg,
+                out,
+                d_l,
+                out,
+                d_l,
+                result.algo.as_ptr(),
+                workspace,
+                ws_bytes,
+                self.stream,
+            );
+            cublasLtMatmulPreferenceDestroy(pref);
+            cublasLtMatrixLayoutDestroy(a_l);
+            cublasLtMatrixLayoutDestroy(b_l);
+            cublasLtMatrixLayoutDestroy(d_l);
+            cublasLtMatmulDescDestroy(desc);
+            if mm != CUBLAS_STATUS_SUCCESS {
+                return Err(MetalTileError::Dispatch(format!(
+                    "cublasLtMatmul(fp4) failed: status {mm} (m={m} n={n} k={k})"
+                )));
+            }
+        }
+        Ok(())
+    }
+
+    /// FP8 (e4m3) GEMM via cuBLASLt with per-tensor f32 DEVICE scale
+    /// pointers (dequant scales; D = scaleX*scaleW * X·Wᵀ). TN layout —
+    /// the FP8-required form. ~2x the f16 rate on GB10, far gentler
+    /// quantization than FP4 (the vendor recipe for this model family runs
+    /// the Mamba/shared/o-proj GEMMs at FP8).
+    #[allow(clippy::too_many_arguments)]
+    pub fn gemm_cublaslt_fp8(
+        &self,
+        x: CUdeviceptr,       // [m, k] e4m3 row-major activation
+        x_scale: CUdeviceptr, // f32 scalar (device)
+        w: CUdeviceptr,       // [n, k] e4m3 row-major weight
+        w_scale: CUdeviceptr, // f32 scalar (device)
+        out: CUdeviceptr,     // [m, n] row-major result
+        m: usize,
+        n: usize,
+        k: usize,
+        out_f32: bool,
+    ) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        let (lt, workspace) = self.cublaslt_ctx()?;
+        let alpha: f32 = 1.0;
+        let beta: f32 = 0.0;
+        unsafe {
+            let mut desc: cublasLtMatmulDesc_t = ptr::null_mut();
+            let s = cublasLtMatmulDescCreate(&mut desc, CUBLAS_COMPUTE_32F, CUDA_R_32F);
+            if s != CUBLAS_STATUS_SUCCESS {
+                return Err(MetalTileError::Dispatch(format!("cublasLtMatmulDescCreate: {s}")));
+            }
+            let opt = CUBLAS_OP_T;
+            let opn = CUBLAS_OP_N;
+            cublasLtMatmulDescSetAttribute(
+                desc,
+                CUBLASLT_MATMUL_DESC_TRANSA,
+                &opt as *const c_int as *const c_void,
+                std::mem::size_of::<c_int>(),
+            );
+            cublasLtMatmulDescSetAttribute(
+                desc,
+                CUBLASLT_MATMUL_DESC_TRANSB,
+                &opn as *const c_int as *const c_void,
+                std::mem::size_of::<c_int>(),
+            );
+            // Per-tensor f32 scales (default SCALAR mode). ptr==0 => skip
+            // (caller folds per-channel/per-token scales as a post-pass).
+            // With the col-major reorder A=W and B=X.
+            if w_scale != 0 {
+                cublasLtMatmulDescSetAttribute(
+                    desc,
+                    CUBLASLT_MATMUL_DESC_A_SCALE_POINTER,
+                    &w_scale as *const CUdeviceptr as *const c_void,
+                    std::mem::size_of::<CUdeviceptr>(),
+                );
+            }
+            if x_scale != 0 {
+                cublasLtMatmulDescSetAttribute(
+                    desc,
+                    CUBLASLT_MATMUL_DESC_B_SCALE_POINTER,
+                    &x_scale as *const CUdeviceptr as *const c_void,
+                    std::mem::size_of::<CUdeviceptr>(),
+                );
+            }
+            let mut a_l: cublasLtMatrixLayout_t = ptr::null_mut();
+            let mut b_l: cublasLtMatrixLayout_t = ptr::null_mut();
+            let mut d_l: cublasLtMatrixLayout_t = ptr::null_mut();
+            cublasLtMatrixLayoutCreate(&mut a_l, CUDA_R_8F_E4M3, k as u64, n as u64, k as i64);
+            cublasLtMatrixLayoutCreate(&mut b_l, CUDA_R_8F_E4M3, k as u64, m as u64, k as i64);
+            cublasLtMatrixLayoutCreate(
+                &mut d_l,
+                if out_f32 { CUDA_R_32F } else { CUDA_R_16F },
+                n as u64,
+                m as u64,
+                n as i64,
+            );
+
+            let mut pref: cublasLtMatmulPreference_t = ptr::null_mut();
+            cublasLtMatmulPreferenceCreate(&mut pref);
+            let ws_bytes: usize = CUBLASLT_WORKSPACE_BYTES;
+            cublasLtMatmulPreferenceSetAttribute(
+                pref,
+                CUBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES,
+                &ws_bytes as *const usize as *const c_void,
+                std::mem::size_of::<usize>(),
+            );
+            let red_mask: u32 = CUBLASLT_REDUCTION_SCHEME_NONE;
+            cublasLtMatmulPreferenceSetAttribute(
+                pref,
+                CUBLASLT_MATMUL_PREF_REDUCTION_SCHEME_MASK,
+                &red_mask as *const u32 as *const c_void,
+                std::mem::size_of::<u32>(),
+            );
+            let mut result = cublasLtMatmulHeuristicResult_t::default();
+            let mut returned: c_int = 0;
+            let hs = cublasLtMatmulAlgoGetHeuristic(
+                lt,
+                desc,
+                a_l,
+                b_l,
+                d_l,
+                d_l,
+                pref,
+                1,
+                &mut result,
+                &mut returned,
+            );
+            if hs != CUBLAS_STATUS_SUCCESS || returned < 1 {
+                cublasLtMatmulPreferenceDestroy(pref);
+                cublasLtMatrixLayoutDestroy(a_l);
+                cublasLtMatrixLayoutDestroy(b_l);
+                cublasLtMatrixLayoutDestroy(d_l);
+                cublasLtMatmulDescDestroy(desc);
+                return Err(MetalTileError::Dispatch(format!(
+                    "cublasLt-fp8: no algo (m={m} n={n} k={k} status={hs} returned={returned})"
+                )));
+            }
+            let mm = cublasLtMatmul(
+                lt,
+                desc,
+                &alpha as *const f32 as *const c_void,
+                w,
+                a_l,
+                x,
+                b_l,
+                &beta as *const f32 as *const c_void,
+                out,
+                d_l,
+                out,
+                d_l,
+                result.algo.as_ptr(),
+                workspace,
+                ws_bytes,
+                self.stream,
+            );
+            cublasLtMatmulPreferenceDestroy(pref);
+            cublasLtMatrixLayoutDestroy(a_l);
+            cublasLtMatrixLayoutDestroy(b_l);
+            cublasLtMatrixLayoutDestroy(d_l);
+            cublasLtMatmulDescDestroy(desc);
+            if mm != CUBLAS_STATUS_SUCCESS {
+                return Err(MetalTileError::Dispatch(format!(
+                    "cublasLtMatmul(fp8) failed: status {mm} (m={m} n={n} k={k})"
+                )));
+            }
+        }
+        Ok(())
+    }
+
+    /// Zero-fill `len` bytes at `ptr`, stream-ordered on the device stream.
+    /// No host staging buffer and no stream drain — the cheap way to seed an
+    /// accumulator (vs uploading a host zero buffer, which for a [s,hid] f32
+    /// accumulator is a multi-MB pageable H2D copy per call).
+    pub fn memset_zero_raw(&self, ptr: CUdeviceptr, len: usize) -> Result<(), MetalTileError> {
+        self.ensure_current();
+        if ptr == 0 || len == 0 {
+            return Ok(());
+        }
+        cu_check(unsafe { cuMemsetD8Async(ptr, 0, len, self.stream) }, "cuMemsetD8Async(zero)")
+    }
+}
diff --git a/crates/metaltile-std/src/ffai/ssm.rs b/crates/metaltile-std/src/ffai/ssm.rs
index e76e2a6e..e4732e63 100644
--- a/crates/metaltile-std/src/ffai/ssm.rs
+++ b/crates/metaltile-std/src/ffai/ssm.rs
@@ -77,7 +77,12 @@ pub fn conv1d_causal_step<T>(
         let next = load(state[(k + 1u32) * n_channels + d]);
         store(state[k * n_channels + d], next);
     }
-    store(state[(kernel_size - 2u32) * n_channels + d], load(x[d]));
+    // Same `kernel_size < 2` hazard as above, but for the tail STORE: the
+    // slot index would wrap to ~4e9 AND `state` has K-1 = 0 slots, so there
+    // is nothing valid to clamp to — skip the store entirely.
+    if kernel_size > 1u32 {
+        store(state[(kernel_size - 2u32) * n_channels + d], load(x[d]));
+    }
 }
 
 // ── Mamba 2 batched-prefill causal depthwise conv1d ─────────────────────
@@ -1101,6 +1106,53 @@ pub fn gated_group_rmsnorm_batched(
 // Inclusive cumsum of A·dt within each chunk → Lcs. One thread per (chunk,head).
 //   A = -exp(a_log[h]),  Lcs[bh, i] = Σ_{k≤i} A·dt[c*L+k]
 // dt layout [T, H]; lcs layout [nc*H, L]. Grid: [nc*H, 1, 1].
+// Same math as `gated_group_rmsnorm_batched`, but stores the final tensor as
+// f16. This is specifically for the CUDA prefill path where the next op is a
+// tensor-core out_proj GEMM that immediately casts the f32 norm output to f16.
+#[kernel]
+pub fn gated_group_rmsnorm_batched_f16out(
+    y: Tensor<f32>,       // [s * di] flat
+    z: Tensor<f32>,       // [s * di] flat
+    w: Tensor<f32>,       // [di]     norm weights (shared across tokens)
+    mut out: Tensor<f16>, // [s * di]
+    eps_buf: Tensor<f32>, // [1]
+    #[constexpr] gs: u32, // group size (512 for Nemotron)
+    #[constexpr] ng: u32, // number of groups per token (8 for Nemotron)
+) {
+    let tg = program_id::<0>();
+    let grp = tg - (tg / ng) * ng;
+    let ti = tg / ng;
+    let rs = ti * ng * gs + grp * gs;
+    let col = tid * 4u32;
+    let in_bounds = col + 3u32 < gs;
+    let safe_col = select(in_bounds, col, 0u32);
+    let sb = rs + safe_col;
+    let y0 = load(y[sb]);
+    let y1 = load(y[sb + 1u32]);
+    let y2 = load(y[sb + 2u32]);
+    let y3 = load(y[sb + 3u32]);
+    let z0 = load(z[sb]);
+    let z1 = load(z[sb + 1u32]);
+    let z2 = load(z[sb + 2u32]);
+    let z3 = load(z[sb + 3u32]);
+    let g0 = y0 * (z0 / (1.0f32 + exp(0.0f32 - z0)));
+    let g1 = y1 * (z1 / (1.0f32 + exp(0.0f32 - z1)));
+    let g2 = y2 * (z2 / (1.0f32 + exp(0.0f32 - z2)));
+    let g3 = y3 * (z3 / (1.0f32 + exp(0.0f32 - z3)));
+    let raw = g0 * g0 + g1 * g1 + g2 * g2 + g3 * g3;
+    let partial = select(in_bounds, raw, 0.0f32);
+    let ssq = reduce_sum(partial);
+    let eps = load(eps_buf[0]);
+    let rms = rsqrt(ssq / (gs.cast::<f32>()) + eps);
+    if in_bounds {
+        let base = rs + col;
+        store(out[base], (g0 * rms * load(w[grp * gs + col])).cast::<f16>());
+        store(out[base + 1u32], (g1 * rms * load(w[grp * gs + col + 1u32])).cast::<f16>());
+        store(out[base + 2u32], (g2 * rms * load(w[grp * gs + col + 2u32])).cast::<f16>());
+        store(out[base + 3u32], (g3 * rms * load(w[grp * gs + col + 3u32])).cast::<f16>());
+    }
+}
+
 #[kernel]
 pub fn ssd_lcs(
     dt: Tensor<f32>,      // [T, H]
@@ -1120,7 +1172,11 @@ pub fn ssd_lcs(
         let mut acc = 0.0f32;
         for i in range(0u32, l, 1u32) {
             let t = c * l + i;
-            let dtv = select(t < t_total, load(dt[t * n_heads + h]), 0.0f32);
+            // `select` pre-evaluates both arms — clamp the index BEFORE the
+            // load so the tail chunk (t ≥ T) never reads past `dt`.
+            let valid = t < t_total;
+            let t_safe = select(valid, t, 0u32);
+            let dtv = select(valid, load(dt[t_safe * n_heads + h]), 0.0f32);
             acc = acc + a * dtv;
             store(lcs[idx * l + i], acc);
         }
@@ -1218,7 +1274,11 @@ pub fn ssd_mmask(
             let c = bh / n_heads;
             let h = bh - c * n_heads;
             let tj = c * l + j;
-            let dtj = select(tj < t_total, load(dt[tj * n_heads + h]), 0.0f32);
+            // Clamp before the load — `select` pre-evaluates both arms, so an
+            // unclamped tj on the zero-padded tail chunk would read OOB.
+            let valid = tj < t_total;
+            let tj_safe = select(valid, tj, 0u32);
+            let dtj = select(valid, load(dt[tj_safe * n_heads + h]), 0.0f32);
             let decay = exp(load(lcs[bh * l + i]) - load(lcs[bh * l + j]));
             store(m_out[e], load(cb[e]) * decay * dtj);
         }
@@ -1283,19 +1343,22 @@ pub fn ssd_recur(
     let idx = program_id::<0>();
     let tot = n_heads * ds * dh;
     if idx < tot {
+        // `sidx`/`stv` (not `s`/`st`): a local whose name prefixes a tensor
+        // param (`s_chunk`, `sin_t`, `state_*`) trips the DSL codegen
+        // local-elision bug → undeclared identifier at dispatch.
         let p = idx - (idx / dh) * dh; // head_dim
-        let s = (idx / dh) - (idx / (dh * ds)) * ds; // state
+        let sidx = (idx / dh) - (idx / (dh * ds)) * ds; // state
         let h = idx / (dh * ds);
-        let mut st = load(state_in[(h * dh + p) * ds + s]);
+        let mut stv = load(state_in[(h * dh + p) * ds + sidx]);
         for c in range(0u32, nc, 1u32) {
             let bh = c * n_heads + h;
             // emit S_inᵀ for this chunk BEFORE applying it
-            store(sin_t[(bh * dh + p) * ds + s], st);
+            store(sin_t[(bh * dh + p) * ds + sidx], stv);
             let alpha = exp(load(lcs[bh * l + (l - 1u32)]));
-            let sc = load(s_chunk[(bh * ds + s) * dh + p]);
-            st = alpha * st + sc;
+            let sc = load(s_chunk[(bh * ds + sidx) * dh + p]);
+            stv = alpha * stv + sc;
         }
-        store(state_out[(h * dh + p) * ds + s], st);
+        store(state_out[(h * dh + p) * ds + sidx], stv);
     }
 }
 
@@ -1421,9 +1484,12 @@ pub fn ssd_g1_cb(
     if i < l {
         if j < l {
             // mmask epilogue: causal (i<j → 0) · decay · dt[t_j, h].
-            // `select` (not if/else) keeps it in the codegen-portable subset.
+            // `select` (not if/else) keeps it in the codegen-portable subset;
+            // clamp tj before the load (select pre-evaluates both arms).
             let tj = c * l + j;
-            let dtj = select(tj < t_total, load(dt[tj * n_heads + h]), 0.0f32);
+            let dt_ok = tj < t_total;
+            let tj_safe = select(dt_ok, tj, 0u32);
+            let dtj = select(dt_ok, load(dt[tj_safe * n_heads + h]), 0.0f32);
             let decay = exp(load(lcs[bz * l + i]) - load(lcs[bz * l + j]));
             let m = select(i < j, 0.0f32, acc * decay * dtj);
             store(out[(bz * l + i) * l + j], m);
@@ -1460,25 +1526,27 @@ pub fn ssd_g4_cs(
     let mut acc = 0.0f32;
     for k0 in range(0u32, ds, 16u32) {
         if tid < 512u32 {
-            // sin_t tile: col p = tgid_x*32 + s/16, state k = k0 + s%16.
-            let s = tid;
-            let p = tgid_x * 32u32 + s / 16u32;
+            // sin_t tile: col p = tgid_x*32 + sl/16, state k = k0 + sl%16.
+            // `sl` (not `s`): `s` prefixes the `sin_t` param name — DSL
+            // codegen local-elision pitfall.
+            let sl = tid;
+            let p = tgid_x * 32u32 + sl / 16u32;
             let valid = p < dh;
             let p_safe = select(valid, p, 0u32);
-            let kk = k0 + s - (s / 16u32) * 16u32;
+            let kk = k0 + sl - (sl / 16u32) * 16u32;
             let wv = select(valid, load(sin_t[w_base + p_safe * ds + kk]), 0.0f32);
-            threadgroup_store("g4cs_w", s, wv);
+            threadgroup_store("g4cs_w", sl, wv);
         }
         if tid >= 512u32 {
-            // C tile: row i = tgid_y*32 + s/16, state k = k0 + s%16.
-            let s = tid - 512u32;
-            let i = tgid_y * 32u32 + s / 16u32;
+            // C tile: row i = tgid_y*32 + sl/16, state k = k0 + sl%16.
+            let sl = tid - 512u32;
+            let i = tgid_y * 32u32 + sl / 16u32;
             let ti = c * l + i;
             let valid = (i < l) & (ti < t_total);
             let ti_safe = select(valid, ti, 0u32);
-            let kk = k0 + s - (s / 16u32) * 16u32;
+            let kk = k0 + sl - (sl / 16u32) * 16u32;
             let cv = select(valid, load(c_mat[(ti_safe * n_groups + g) * ds + kk]), 0.0f32);
-            threadgroup_store("g4cs_c", s, cv);
+            threadgroup_store("g4cs_c", sl, cv);
         }
         threadgroup_barrier();
         for k in range(0u32, 16u32, 1u32) {
diff --git a/crates/metaltile-std/src/mlx/rms_norm.rs b/crates/metaltile-std/src/mlx/rms_norm.rs
index 06ca47dd..b2a584dd 100644
--- a/crates/metaltile-std/src/mlx/rms_norm.rs
+++ b/crates/metaltile-std/src/mlx/rms_norm.rs
@@ -154,6 +154,41 @@ pub fn mt_rms_norm<T>(
 /// Algorithm-identical to `mt_rms_norm`: f32 accumulator for the
 /// sum-of-squares, threadgroup-wide `reduce_sum`, `rsqrt(ssq/n + eps)`
 /// scaling, per-element output store rounded through `T`.
+/// Row-wise RMSNorm with only an f16 output surface. This is equivalent to
+/// `cast_f32_f16(mt_rms_norm(x, w))`, but avoids materializing the f32 norm
+/// when the next consumer is a tensor-core projection.
+#[kernel]
+pub fn mt_rms_norm_f16out<T>(
+    x: Tensor<T>,
+    w: Tensor<T>,
+    mut out: Tensor<f16>,
+    eps_buf: Tensor<f32>,
+    #[constexpr] n: u32,
+) {
+    let row = program_id::<0>();
+    let rs = row * n;
+    let col = tid * 4u32;
+    let in_bounds = col + 3u32 < n;
+    let safe_col = select(in_bounds, col, 0u32);
+    let safe_base = rs + safe_col;
+    let base = rs + col;
+    let x0 = load(x[safe_base]).cast::<f32>();
+    let x1 = load(x[safe_base + 1u32]).cast::<f32>();
+    let x2 = load(x[safe_base + 2u32]).cast::<f32>();
+    let x3 = load(x[safe_base + 3u32]).cast::<f32>();
+    let raw_ssq = x0 * x0 + x1 * x1 + x2 * x2 + x3 * x3;
+    let partial_ssq = select(in_bounds, raw_ssq, 0.0f32);
+    let tg_ssq = reduce_sum(partial_ssq);
+    let eps = load(eps_buf[0]);
+    let rms = rsqrt(tg_ssq / n + eps);
+    if in_bounds {
+        store(out[base], (x0 * rms * load(w[col]).cast::<f32>()).cast::<f16>());
+        store(out[base + 1u32], (x1 * rms * load(w[col + 1u32]).cast::<f32>()).cast::<f16>());
+        store(out[base + 2u32], (x2 * rms * load(w[col + 2u32]).cast::<f32>()).cast::<f16>());
+        store(out[base + 3u32], (x3 * rms * load(w[col + 3u32]).cast::<f32>()).cast::<f16>());
+    }
+}
+
 #[kernel]
 pub fn mt_rms_norm_small<T>(
     x: Tensor<T>,