feat(runtime): NVFP4 module split, CUBIN opt-in, col-copy guard, f16out kernels

.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',
 ]
 

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;

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"))
+    }
+}