feat: fmt=8 (fp8-e4m3) decode on the kv_b absorb path, CPU and CUDA

c/Makefile.deepseek-v4

@@ -264,20 +264,20 @@ $(REGISTRY_OBJ): expert_store_registry.c expert_store_registry.h expert_store.h
 	$(CC) $(CFLAGS) -c expert_store_registry.c -o $@
 
 $(TARGET_OBJS) $(TEST_UNIT_OBJS): %.o: deepseek_v4.c deepseek_v4.h \
-	deepseek_v4_internal.h deepseek_v4_dspark.inc st.h json.h compat.h tensor.h quant.h \
+	deepseek_v4_internal.h deepseek_v4_dspark.inc st.h json.h compat.h tensor.h quant.h fp8_format.h \
 	route_trace.h \
 	native_quant.h native_quant_batch.h native_quant_dual.h \
 	native_quant_fp4_rows16.h expert_store_registry.h $(V4_FLAGS_STAMP)
 	$(CC) $(CFLAGS) -D$* -c deepseek_v4.c -o $@
 
 $(V4_HOT_TEST_OBJ): deepseek_v4.c deepseek_v4.h deepseek_v4_internal.h \
-		st.h json.h compat.h tensor.h quant.h route_trace.h \
+		st.h json.h compat.h tensor.h quant.h fp8_format.h route_trace.h \
 		native_quant.h native_quant_fp4_rows16.h expert_store_registry.h $(V4_FLAGS_STAMP)
 	$(CC) $(CFLAGS) -DCOLI_V4_TEST_HOOKS \
 		-DCOLI_V4_UNIT_EXPERT_STORE_HOT_ROWS16 -c deepseek_v4.c -o $@
 
 $(V4_BATCH_TEST_OBJ): deepseek_v4.c deepseek_v4.h deepseek_v4_internal.h \
-		tensor.h quant.h native_quant.h native_quant_batch.h $(V4_FLAGS_STAMP)
+		tensor.h quant.h fp8_format.h native_quant.h native_quant_batch.h $(V4_FLAGS_STAMP)
 	$(CC) $(CFLAGS) -DCOLI_V4_TEST_HOOKS \
 		-DCOLI_V4_UNIT_NATIVE_QUANT_BATCH -c deepseek_v4.c -o $@
 

c/backend_cuda.cu

@@ -1,7 +1,12 @@
 #include "backend_cuda.h"
+#include "fp8_format.h"   /* FP8_BLOCK: the shared fmt=8 scale-block edge (see that header) */
 
 #include "backend_gpu_compat.h"
 
+static_assert(FP8_BLOCK == 128, "fmt=8 on-disk containers carry ceil(dim/128)-edged scale "
+              "grids (mint tool, docs/FORMATS.md); FP8_BLOCK is container format, not a "
+              "tunable -- an edit here is a format change");
+
 /* Optional fmt=8 decode candidate (COLI_CUDA_F8_WARP=2): cuda_fp8.h maps
  * __nv_cvt_fp8_to_halfraw to an sm_89+ cvt instruction, with a bit-manip
  * fallback below 890. CUDA-only; the HIP build keeps the LUT decode. */
@@ -287,13 +292,14 @@ __device__ static inline float mx4_weight_at(const uint8_t *q, int i) {
  * branch and the fall-through is a refusal.
  *
  * It used to be the other way round: int2 was the fall-through, so every format
- * this function does not decode -- fmt=5 (int3-g64), fmt=6 (E8/IQ3), fmt=8
- * (fp8-e4m3), and anything added later -- was read as 2-bit values and returned
- * numbers. Meanwhile the CPU functions doing the same job on the same tensor,
- * qt_addrow and qt_matvec_rows (colibri.c), both exit(1) naming the function and
- * the fmt. Two backends, identical unsupported input, one refusing and one
- * fabricating: that asymmetry is the defect, independent of any particular
- * format's arrival.
+ * this function does not decode -- fmt=5 (int3-g64), fmt=6 (E8/IQ3), and
+ * anything added later -- was read as 2-bit values and returned numbers.
+ * (fmt=8 was in that misread set too, then refused, until it gained its own
+ * explicit branch below for the absorb path.) Meanwhile the CPU functions
+ * doing the same job on the same tensor, qt_addrow and qt_matvec_rows
+ * (colibri.c), both exit(1) naming the function and the fmt. Two backends,
+ * identical unsupported input, one refusing and one fabricating: that
+ * asymmetry is the defect, independent of any particular format's arrival.
  *
  * WHY __trap() AND NOT A DIAGNOSTIC. This is device code inside a running
  * kernel; there is no stderr to name the tensor on and no way to unwind. __trap
@@ -322,6 +328,15 @@ __device__ static float weight_at(const void *weights, int fmt, size_t row, int
     const uint8_t *base = static_cast<const uint8_t *>(weights) + row;
     if (fmt == 0) return reinterpret_cast<const float *>(base)[i];
     if (fmt == 1) return static_cast<float>(reinterpret_cast<const int8_t *>(base)[i]);
+    /* fmt=8 (fp8-e4m3): raw byte, same layout as fmt=1 (row_bytes(8,I)==I), decoded
+     * through the shared c_e4m3 LUT (same table quant_matmul's fmt==8 branch reads,
+     * uploaded once by coli_cuda_fp8_set_lut). Callers gate on the LUT being live
+     * before a fmt=8 tensor ever reaches this function (coli_cuda_tensor_upload
+     * refuses the upload otherwise), so the table is always populated here. Returns
+     * the decoded WEIGHT only, unscaled -- absorb_scale below applies the
+     * per-128x128-block scale, exactly like every other quantized fmt returns
+     * unscaled through this function. */
+    if (fmt == 8) return c_e4m3[base[i]];
     const uint8_t *q = base;
     if (fmt == 2 || fmt == 4) {                               /* fmt=4: same nibble layout */
         uint8_t v = q[i >> 1];
@@ -335,13 +350,34 @@ __device__ static float weight_at(const void *weights, int fmt, size_t row, int
     return 0.0f;   /* not reached: __trap() does not return */
 }
 
-/* Scale for output `row`, input element `k`. fmt=4 (grouped int4) stores ng
- * scales per row at scales[row*ng + k/gs]; every other quantized format has
- * one scale per row at scales[row]. Mirrors quant_matmul's fmt==4 branch so the
+/* Scale for output `row`, input element `k`. Three layouts reach this: fmt=4
+ * (grouped int4) stores ng scales per row at scales[row*ng + k/gs]; fmt=8
+ * (fp8-e4m3) stores one scale per 128x128 BLOCK, block-row-major, and is
+ * handled by its own branch below; every OTHER quantized format has one scale
+ * per row at scales[row]. Mirrors quant_matmul's fmt==4 branch so the
  * attention absorb kernels apply per-group scales instead of the per-row
  * (fmt=2) semantic that crashed #298's g64 kv_b. */
 __device__ static float absorb_scale(const float *wscale, int fmt, int gs, int ng, int row, int k) {
     if (!fmt) return 1.f;
+    if (fmt == 8) {
+        /* fp8-e4m3: one f32 scale per 128x128 BLOCK, block-row-major
+         * ([ceil(O/128), ceil(I/128)]), exactly quant_matmul's fmt==8 indexing
+         * (scl[i >> 7] on a scale row selected by o >> 7) and matmul_fp8's CPU
+         * reference (quant.h). `ng` here is coli_cuda_tensor_upload's t->ng,
+         * which for fmt=8 is set to ceil(I/128) specifically (not the fmt=4
+         * group count) -- see the upload-time assignment there. `gs` is unused
+         * for fmt=8 (always 0, only fmt=4 sets it), so the block edge is the
+         * fixed FP8_BLOCK constant (fp8_format.h, shared with the CPU side),
+         * not a caller-supplied group size. Rounding note: the GEOMETRY here
+         * matches quant_matmul_f8w/matmul_fp8, but their fp8 accumulation
+         * convention (f32 partial per block, scale once per partial, double
+         * across blocks) is NOT carried into the absorb kernels -- they apply
+         * the scale per element into a float accumulator, matching their own
+         * fmt=4 arm's long-standing behavior; CPU-vs-CUDA absorb divergence
+         * is an accepted, documented class (#510). */
+        int rowBlk = row / FP8_BLOCK, colBlk = k / FP8_BLOCK;
+        return wscale[(size_t)rowBlk * ng + colBlk];
+    }
     if (fmt != 4) return wscale[row];
     int g = k / gs; if (g >= ng) g = ng - 1;   /* tail of the last (partial) group */
     return wscale[(size_t)row * ng + g];
@@ -558,9 +594,9 @@ __global__ static void quant_matmul(float *y, const float *x, const void *weight
          * the ORIGINAL dense path, kept for COLI_CUDA_F8_WARP=0; the default
          * routes fmt=8 to quant_matmul_f8w instead (quant_matmul_launch). */
         const uint8_t *wrow = static_cast<const uint8_t *>(weights) + row;
-        const float *scl = scales + (size_t)(o >> 7) * (size_t)((I + 127) >> 7);
+        const float *scl = scales + (size_t)(o / FP8_BLOCK) * (size_t)((I + FP8_BLOCK - 1) / FP8_BLOCK);
         for (int i = threadIdx.x; i < I; i += blockDim.x)
-            sum += xs[i] * c_e4m3[wrow[i]] * scl[i >> 7];
+            sum += xs[i] * c_e4m3[wrow[i]] * scl[i / FP8_BLOCK];
     } else {
         for (int i = threadIdx.x; i < I; i += blockDim.x)
             sum += xs[i] * weight_at(weights, fmt, row, i);
@@ -1272,10 +1308,15 @@ extern "C" int coli_cuda_init(const int *devices, int count) {
         }
     }
     if (g_nctx) {
-        int same = count == g_nctx;
-        for (int i = 0; same && i < count; i++) same = devices[i] == g_ctx[i].device;
-        if (!same) std::fprintf(stderr, "[CUDA] device list change requires shutdown first\n");
-        return same;
+        /* Same decision as before, routed through the shared predicate in
+         * backend_cuda.h so a host-side test can pin it without nvcc; the
+         * return value is unchanged (1 for the same set, 0 otherwise). */
+        int live[COLI_CUDA_MAX_DEVICES];
+        for (int i = 0; i < g_nctx; i++) live[i] = g_ctx[i].device;
+        int d = coli_cuda_init_disposition(g_nctx, count, devices, live);
+        if (d == COLI_CUDA_INIT_REFUSE)
+            std::fprintf(stderr, "[CUDA] device list change requires shutdown first\n");
+        return d == COLI_CUDA_INIT_ACCEPT;
     }
     for (int i = 0; i < count; i++) {
         int device = devices[i];
@@ -1356,6 +1397,22 @@ extern "C" void coli_cuda_shutdown(void) {
         ctx->group_desc=nullptr; ctx->group_desc_cap=0;
     }
     g_nctx = 0;
+    /* g_fp8_lut_ready is PROCESS-WIDE while the e4m3 table (c_e4m3, a
+     * __constant__ device symbol whose lifetime is the CUDA primary context,
+     * not this file's host-side DeviceContext structs) is PER-DEVICE. A later
+     * coli_cuda_init may select a device the previous span never published to;
+     * without this reset the upload gate (g_fp8_lut_ready, checked in
+     * coli_cuda_tensor_upload) would still be satisfied from the PREVIOUS
+     * boot and admit fmt=8 tensors whose kernels there decode against an
+     * unwritten (zero) table: silent all-zero weights, the exact
+     * fabricated-numbers failure mode the format gates exist to refuse.
+     * Reset so every boot must publish its own LUT (coli_cuda_fp8_set_lut)
+     * before any fmt=8 upload. Shutdown is the ONLY site that needs to clear
+     * the flag: coli_cuda_init refuses a re-init that names a different
+     * device set (it returns early while g_nctx is non-zero, leaving the
+     * existing contexts and their published table untouched), so the device
+     * set can only WIDEN by passing through here first. */
+    g_fp8_lut_ready = 0;
 #ifdef COLI_ANS
     if(g_ans_sidecar){std::fclose(g_ans_sidecar);g_ans_sidecar=nullptr;}
 #if defined(__linux__)
@@ -1444,7 +1501,9 @@ extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
     /* fmt=6 keeps its scales inside each 98-byte block, so it is the one
      * quantized format that legitimately arrives with scales == NULL. */
     if (!rb || (fmt && fmt != 6 && !scales)) return 0;
-    if (fmt == 8 && !g_fp8_lut_ready) return 0;   /* kernels would read a zero LUT */
+    /* kernels would read a zero LUT; shared predicate, pinned by
+     * tests/test_cuda_lut_gate.c without a CUDA toolchain */
+    if (!coli_cuda_fp8_gate_admits(fmt, g_fp8_lut_ready)) return 0;
     ColiCudaTensor *t = static_cast<ColiCudaTensor *>(std::calloc(1, sizeof(*t)));
     if (!t) return 0;
     t->fmt = fmt; t->I = I; t->O = O; t->device = device; t->weight_bytes = rb * (size_t)O;
@@ -1455,9 +1514,9 @@ extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
         t->ng = (I + 31) / 32;
         t->scale_count = (size_t)O * t->ng;
     }
-    if (fmt == 8) {   /* per-128x128-block scales: [ceil(O/128), ceil(I/128)] */
-        t->ng = (I + 127) / 128;
-        t->scale_count = (size_t)((O + 127) / 128) * (size_t)t->ng;
+    if (fmt == 8) {   /* per-block scales: [ceil(O/FP8_BLOCK), ceil(I/FP8_BLOCK)] (fp8_format.h) */
+        t->ng = (int)fp8_nblk(I);
+        t->scale_count = (size_t)fp8_nblk(O) * (size_t)t->ng;
     }
     if (!cuda_ok(cudaMalloc(&t->weights, t->weight_bytes), "tensor allocation")) {
         coli_cuda_tensor_free(t);
@@ -2278,18 +2337,20 @@ extern "C" const float *coli_cuda_expert_group_take(int device) {
 
 
 /* The absorb kernels decode `w` through weight_at + absorb_scale, which know
- * per-row and fmt=4 group scales only. Refuse anything else (fmt=5/6/8) rather
- * than mis-decode it — the caller keeps its CPU attention path. (`proj`
- * tensors are exempt: they run through quant_matmul, which dispatches every
- * format it uploads.) A dedicated block-scale absorb for fmt=8 is follow-up
- * work, same shape as routing fmt=4 through the grouped kernels was.
+ * per-row scales, fmt=4 group scales, and fmt=8 per-128x128-block scales.
+ * Refuse anything else (fmt=5/6/7) rather than mis-decode it — the caller
+ * keeps its CPU attention path. (`proj` tensors are exempt: they run through
+ * quant_matmul, which dispatches every format it uploads.) fmt=8 support
+ * funnels through this one predicate for all the absorb host wrappers below,
+ * so none of them needed a separate change.
  *
  * The admissible set is weight_at's own, taken from the shared predicate rather
- * than restated as `fmt <= 4`: this gate and weight_at's device-side backstop
- * must not be able to drift apart, and the old inequality also admitted
- * NEGATIVE fmt values, which weight_at would then have fallen through on. Same
- * truth table for every fmt a container can actually carry (0..8), so no
- * existing container changes behaviour here. */
+ * than restated as an inequality: this gate and weight_at's device-side
+ * backstop must not be able to drift apart, and the old `fmt <= 4` also
+ * admitted NEGATIVE fmt values, which weight_at would then have fallen through
+ * on. A fmt=8 tensor implies a live e4m3 LUT (upload refuses it otherwise --
+ * see the predicate's caveat note in backend_cuda.h), so no extra gate is
+ * needed here. */
 static int absorb_fmt_ok(const ColiCudaTensor *w){
     return w && coli_cuda_weight_at_supported(w->fmt);
 }

c/backend_cuda.h

@@ -23,7 +23,9 @@ extern "C" {
 
 /* Weight formats the generic per-element device decoder (weight_at,
  * backend_cuda.cu) can actually decode: f32, int8-row, int4 nibbles (fmt=2 and
- * the grouped fmt=4, same packing), and int2. Nothing else.
+ * the grouped fmt=4, same packing), int2, and fmt=8 (fp8-e4m3 raw bytes,
+ * decoded through the c_e4m3 LUT -- absorb-path support; absorb_scale supplies
+ * its per-128x128-block scale). Nothing else.
  *
  * WHY THIS IS A PREDICATE AND NOT A COMMENT. weight_at used to END in the int2
  * decode as an unguarded fall-through, so ANY other format handed to it -- a
@@ -42,12 +44,60 @@ extern "C" {
  * same arrangement colibri.c uses for metal_fused_fmt_ok.
  *
  * NOT a statement about which formats the CUDA BACKEND supports: quant_matmul
- * has its own explicit branches for fmt=6 (E8/IQ3), fmt=7 (MXFP4) and fmt=8
- * (fp8-e4m3) that never route through weight_at. This predicate is scoped to
- * weight_at's own dispatch, which is what the absorb and grouped-expert kernels
- * decode through. */
+ * has its own explicit branches for fmt=6 (E8/IQ3) and fmt=7 (MXFP4) that
+ * never route through weight_at (and its own fmt=8 branch for the dense path
+ * -- weight_at's fmt=8 branch serves the absorb kernels, which share the same
+ * c_e4m3 LUT). This predicate is scoped to weight_at's own dispatch, which is
+ * what the absorb and grouped-expert kernels decode through.
+ *
+ * fmt=8 CAVEAT, stated because the truth table alone cannot carry it: a fmt=8
+ * decode additionally requires the e4m3 LUT to have been published to the
+ * configured devices (coli_cuda_fp8_set_lut). The exact mechanism, so the
+ * claim cannot outrun it: coli_cuda_fp8_set_lut copies the table into every
+ * context live AT CALL TIME and sets a process-wide flag; the flag gates
+ * fmt=8 uploads (coli_cuda_tensor_upload refuses until it is set).
+ * coli_cuda_shutdown clears the flag; coli_cuda_init never writes it. That is
+ * enough because init will not rebuild contexts underneath a live set: a
+ * re-init naming the same device set returns success and leaves the contexts,
+ * and the table published to them, untouched, while one naming a different
+ * device set is refused before any context is rebuilt. So the device set
+ * cannot widen past what the last publish covered without going through
+ * shutdown, and no fmt=8 ColiCudaTensor can reach a kernel whose device has
+ * an unwritten table. This predicate
+ * deliberately does not restate that gate: it answers "does weight_at have a
+ * decode branch for this fmt", which is the question the launch-site gates
+ * and the device-side __trap() backstop share. */
 static inline int coli_cuda_weight_at_supported(int fmt) {
-    return fmt == 0 || fmt == 1 || fmt == 2 || fmt == 3 || fmt == 4;
+    return fmt == 0 || fmt == 1 || fmt == 2 || fmt == 3 || fmt == 4 || fmt == 8;
+}
+
+/* The two decisions the fmt=8 LUT gate rests on, as pure predicates. They live
+ * here rather than inline in backend_cuda.cu so a host-side test can pin them
+ * with no CUDA toolchain and no GPU (tests/test_cuda_lut_gate.c). backend_cuda.cu
+ * calls BOTH at the real decision sites, so the test pins the engine's own
+ * logic rather than a second copy that could drift from it -- which is the
+ * failure this factoring exists to prevent, the gate having no CI reach
+ * otherwise. */
+
+/* Does the upload gate admit this tensor? Only fmt=8 needs the published
+ * table; every other format decodes without one. */
+static inline int coli_cuda_fp8_gate_admits(int fmt, int lut_ready) {
+    return fmt != 8 || lut_ready != 0;
+}
+
+/* What coli_cuda_init must do with a request while a device set may be live.
+ * BUILD: nothing is live, build the contexts. ACCEPT: the same set is already
+ * live -- return success and touch nothing, so the table published to those
+ * contexts stays valid. REFUSE: a different set is live -- refuse before
+ * rebuilding anything, so the set cannot widen past the last publish. */
+enum { COLI_CUDA_INIT_BUILD = 0, COLI_CUDA_INIT_ACCEPT = 1, COLI_CUDA_INIT_REFUSE = -1 };
+static inline int coli_cuda_init_disposition(int nctx, int count,
+                                             const int *want, const int *live) {
+    int i;
+    if (nctx <= 0) return COLI_CUDA_INIT_BUILD;
+    if (count != nctx) return COLI_CUDA_INIT_REFUSE;
+    for (i = 0; i < count; i++) if (want[i] != live[i]) return COLI_CUDA_INIT_REFUSE;
+    return COLI_CUDA_INIT_ACCEPT;
 }
 
 /* Opaque, persistent device copy of one resident quantized tensor. */

c/backend_metal.mm

@@ -767,7 +767,7 @@ static size_t fmt_bytes(int fmt, int I, int O) {
 // Grouped-int4 (fmt=4) scale-array size: one f32 per gsz-element group, per row -> O*ceil(I/gsz).
 // fp8 (fmt=8) scale-array size: one f32 per 128x128 BLOCK -> ceil(O/128)*ceil(I/128) (2D,
 // not per-row -- quant.h isn't included here, so the ceil-div is inlined rather than sharing
-// colibri.c's qt_scale_bytes/quant.h's fp8_nblk). The block is a fixed 128x128, so gs is
+// colibri.c's qt_scale_bytes/fp8_format.h's fp8_nblk). The block is a fixed 128x128, so gs is
 // ignored for fmt==8. f32 is this build's implemented scale
 // ENCODING for fmt=8 (see quant.h/colibri.c) -- this file has no reason to know that a
 // UE8M0 encoding exists at all: qt_resolve_fmt refuses it on the CPU read path before any