From f1f7f9abd0ebc20929cea20d1817426183584ba9 Mon Sep 17 00:00:00 2001 From: akol1 <15747200+akol1@users.noreply.github.com> Date: Thu, 15 Oct 2026 21:08:12 -0700 Subject: [PATCH 1/4] cuda: serve losslessly packed MXFP8 projections through native qlinear --- zig/kernels/cuda/qmmf.cu | 17 +++- zig/src/cuda/mxfp8.md | 27 +++++++ zig/src/cuda/mxfp8.zig | 41 ++++++++++ zig/src/cuda/qlinear.zig | 9 ++- zig/src/cuda/qmmf.zig | 14 ++-- zig/src/cuda/root.zig | 2 + zig/tests/cuda/main.zig | 3 + zig/tests/cuda/mxfp8_tests.zig | 143 +++++++++++++++++++++++++++++++++ 8 files changed, 247 insertions(+), 9 deletions(-) create mode 100644 zig/src/cuda/mxfp8.md create mode 100644 zig/src/cuda/mxfp8.zig create mode 100644 zig/tests/cuda/mxfp8_tests.zig diff --git a/zig/kernels/cuda/qmmf.cu b/zig/kernels/cuda/qmmf.cu index 4703273410..027b8ef630 100644 --- a/zig/kernels/cuda/qmmf.cu +++ b/zig/kernels/cuda/qmmf.cu @@ -32,6 +32,11 @@ __device__ __forceinline__ uint32_t fp8pair(uint32_t w) { return r; } +// UE8M0 has no zero encoding: byte 0 is 2^-127, byte 255 is NaN. +__device__ __forceinline__ float e8m0f(uint8_t b) { + return __uint_as_float(b == 0 ? 0x00400000u : b == 255 ? 0x7FC00000u : uint32_t(b) << 23); +} + __device__ __forceinline__ uint32_t comp(const uint4& v, int c) { return c == 0 ? v.x : c == 1 ? v.y : c == 2 ? v.z : v.w; } __device__ __forceinline__ float e4m3f(uint8_t b) { @@ -193,8 +198,8 @@ __global__ void __launch_bounds__(WM * WN * 32) qmmf_kernel( s0 = e4m3f(static_cast(sq[j][0] >> (8 * blk))); s1 = e4m3f(static_cast(sq[j][1] >> (8 * blk))); } else { - s0 = __int_as_float(static_cast((sq[j][0] >> (8 * blk)) & 0xFFu) << 23); - s1 = __int_as_float(static_cast((sq[j][1] >> (8 * blk)) & 0xFFu) << 23); + s0 = e8m0f(static_cast(sq[j][0] >> (8 * blk))); + s1 = e8m0f(static_cast(sq[j][1] >> (8 * blk))); } #pragma unroll for (int i = 0; i < T::MT; ++i) @@ -293,6 +298,14 @@ __global__ void reduce_kernel(const float* __restrict__ part, void* __restrict__ else reinterpret_cast<__nv_bfloat16*>(out)[i] = __float2bfloat16_rn(acc); } +template __global__ void qmmf_kernel(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int); +template __global__ void qmmf_kernel(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int); +template __global__ void qmmf_kernel(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int); +template __global__ void qmmf_kernel(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int); +template __global__ void qmmf_kernel(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int); +template __global__ void qmmf_kernel(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int); +template __global__ void qmmf_kernel(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int); + // The instances the Zig launchers load by name (the Python host instantiates these through its launch switch). template __global__ void qmmf_kernel(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int); template __global__ void qmmf_kernel(const __nv_bfloat16* __restrict__, const unsigned char* __restrict__, const uint8_t* __restrict__, float, void* __restrict__, float* __restrict__, int, int, int, int, int, int, int); diff --git a/zig/src/cuda/mxfp8.md b/zig/src/cuda/mxfp8.md new file mode 100644 index 0000000000..5878df43eb --- /dev/null +++ b/zig/src/cuda/mxfp8.md @@ -0,0 +1,27 @@ +# Native MXFP8 projections + +`cuda.mxfp8` packs row-major E4M3 code bytes and one UE8M0 scale byte per +32 input channels for `qlinear.Weight.mxfp8`. Both prompt and decode calls use +the lane matmul, with BF16 inputs, FP32 accumulation and BF16 output. Packing +permutes bytes; it does not requantize weights or expand them to BF16. + +The logical shape is `[N, K]`, with positive dimensions and K divisible by 64. +`packCodes` writes `padded(N) * K` bytes; `packScales` writes +`padded(N) * K/32` bytes. Rows are padded to 128; padded codes are zero and +padded scales are one. `weight` returns the borrowed device view. The caller +owns its device buffers and the lane's split-K scratch. + +UE8M0 bytes 0 through 254 mean `2^(byte - 127)`; byte 255 is NaN. In +particular, byte zero is a nonzero subnormal FP32 scale. See the +[CUDA datatype contract](https://docs.nvidia.com/cuda/archive/12.8.0/cuda-math-api/cuda_math_api/struct____nv__fp8__e8m0.html). + +`tf-cuda-test mxfp8` checks all 256 scale encodings through the production +matmul, then compares synthetic projections with an independent FP64 +calculation. It also checks strided activations, partial output tiles, +row counts 1/7/17/33/256, prompt/decode dispatch, and cluster versus explicit +split reduction. Host tests cover scale permutation, padding, shared-memory +sizes, symbols and format refusals. + +This primitive does not admit new checkpoint families or certify model +correctness, tensor parallelism or performance. Rebuild the qmmf fatbin with +the native runtime; previous fatbins lack the MXFP8 instantiations. diff --git a/zig/src/cuda/mxfp8.zig b/zig/src/cuda/mxfp8.zig new file mode 100644 index 0000000000..7c57aa50df --- /dev/null +++ b/zig/src/cuda/mxfp8.zig @@ -0,0 +1,41 @@ +//! E4M3 codes with one UE8M0 byte per 32 inputs, repacked losslessly for the lane matmul. +const std = @import("std"); +const fp8 = @import("fp8.zig"); +const qmmf = @import("qmmf.zig"); + +pub const group = 32; +pub const padded = fp8.padded; +pub const packCodes = fp8.packCodes; + +/// Row-major [n, k/32] bytes to [npad/64][k/64][64][2]; padded rows have scale 1. +pub fn packScales(out: []u8, scales: []const u8, n: usize, k: usize) void { + const npad = padded(n); + const kg = k / qmmf.group; + std.debug.assert(k % qmmf.group == 0 and out.len == npad * (k / group) and scales.len == n * (k / group)); + for (0..npad / 64) |t| for (0..kg) |g| for (0..64) |r| for (0..2) |h| { + const row = t * 64 + r; + out[((t * kg + g) * 64 + r) * 2 + h] = if (row < n) scales[row * (k / group) + g * 2 + h] else 127; + }; +} + +pub fn weight(codes: u64, scales: u64, n: usize, k: usize) qmmf.Weight { + return .{ .mode = .mxfp8, .codes = codes, .scales = scales, .n = @intCast(n), .k = @intCast(k), .npad = @intCast(padded(n)) }; +} + +test "every scale byte lands in its independent 32-input group, including padded rows" { + const a = std.testing.allocator; + const n = 130; + const k = 256; + const scales = try a.alloc(u8, n * (k / group)); + defer a.free(scales); + for (scales, 0..) |*v, i| v.* = @truncate(i * 13 + 7); + const out = try a.alloc(u8, padded(n) * (k / group)); + defer a.free(out); + packScales(out, scales, n, k); + for (0..padded(n)) |row| for (0..k / group) |g| { + const at = (((row / 64) * (k / 64) + g / 2) * 64 + row % 64) * 2 + g % 2; + try std.testing.expectEqual(if (row < n) scales[row * (k / group) + g] else @as(u8, 127), out[at]); + }; + const w = weight(1, 2, n, k); + try std.testing.expect(w.mode == .mxfp8 and w.npad == 256 and w.scale == 1.0); +} diff --git a/zig/src/cuda/qlinear.zig b/zig/src/cuda/qlinear.zig index 6d6a56f5a3..b924132f79 100644 --- a/zig/src/cuda/qlinear.zig +++ b/zig/src/cuda/qlinear.zig @@ -33,6 +33,7 @@ pub const Affine4 = struct { pub const Weight = union(enum) { affine4: Affine4, fp8g: qmmf.Weight, + mxfp8: qmmf.Weight, // e4m3 with UE8M0 scale bytes per 32 inputs nvfp4: qmmf.Weight, fp8: qmmf.Weight, // e4m3 with one fp32 scale (ModelOpt FP8) bf16: qmmf.Weight, // unquantized: the checkpoint's bf16 values @@ -40,14 +41,14 @@ pub const Weight = union(enum) { pub fn outputs(w: Weight) usize { return switch (w) { .affine4 => |q| q.n, - .fp8g, .nvfp4, .fp8, .bf16 => |q| q.n, + .fp8g, .mxfp8, .nvfp4, .fp8, .bf16 => |q| q.n, }; } pub fn inputs(w: Weight) usize { return switch (w) { .affine4 => |q| q.k, - .fp8g, .nvfp4, .fp8, .bf16 => |q| q.k, + .fp8g, .mxfp8, .nvfp4, .fp8, .bf16 => |q| q.k, }; } }; @@ -70,6 +71,7 @@ pub const Linear = struct { return a.decode(s, in.x, in.sums, q, out, rows); }, .fp8g => |q| return l.onLane(s, q, .fp8g, in, out, rows), + .mxfp8 => |q| return l.onLane(s, q, .mxfp8, in, out, rows), .nvfp4 => |q| return l.onLane(s, q, .fp4, in, out, rows), .fp8 => |q| return l.onLane(s, q, .fp8, in, out, rows), .bf16 => |q| return l.onLane(s, q, .bf16, in, out, rows), @@ -85,6 +87,7 @@ pub const Linear = struct { return a.prompt(s, in.x, q, out, rows); }, .fp8g => |q| return l.onLane(s, q, .fp8g, in, out, rows), + .mxfp8 => |q| return l.onLane(s, q, .mxfp8, in, out, rows), .nvfp4 => |q| return l.onPrompt(s, q, .fp4, in, out, rows), .fp8 => |q| return l.onPrompt(s, q, .fp8, in, out, rows), .bf16 => |q| return l.onLane(s, q, .bf16, in, out, rows), @@ -330,5 +333,7 @@ test "a lane weight whose mode is not its format's is refused before any launch" const fp4: qmmf.Weight = .{ .mode = .fp4, .codes = 0, .scales = 0, .n = 64, .k = 64, .npad = 128 }; const s: Stream = undefined; try std.testing.expectError(error.FormatMismatch, l.decode(s, .{ .fp8g = fp4 }, .{ .x = 0 }, 0, 1)); + try std.testing.expectError(error.FormatMismatch, l.decode(s, .{ .mxfp8 = fp4 }, .{ .x = 0 }, 0, 1)); + try std.testing.expectError(error.FormatMismatch, l.prompt(s, .{ .mxfp8 = fp4 }, .{ .x = 0 }, 0, 1)); try std.testing.expectError(error.FormatNotLoaded, l.prompt(s, .{ .nvfp4 = fp4 }, .{ .x = 0 }, 0, 1)); } diff --git a/zig/src/cuda/qmmf.zig b/zig/src/cuda/qmmf.zig index 054e2c34bb..3e2b4168a7 100644 --- a/zig/src/cuda/qmmf.zig +++ b/zig/src/cuda/qmmf.zig @@ -1,4 +1,4 @@ -//! The lane matmul (qmmf.cu) for FP4, FP8, FP8G and BF16 weights, with its kernel, tile and slices. +//! The lane matmul (qmmf.cu) for FP4, FP8, MXFP8, FP8G and BF16 weights, with its kernel, tile and slices. const std = @import("std"); const driver = @import("driver.zig"); const module = @import("module.zig"); @@ -17,9 +17,9 @@ const bn = 64; const threads = 128; // WM 1 x WN 4 warps /// The packed formats (qmmf.cu Mode): FP4 with e4m3 scales a 16, e4m3 with one scale or fp32 a 64, and bf16. -pub const Mode = enum(u3) { fp4 = 0, fp8 = 1, fp8g = 3, bf16 = 4 }; +pub const Mode = enum(u3) { fp4 = 0, fp8 = 1, mxfp8 = 2, fp8g = 3, bf16 = 4 }; -const modes = [_]Mode{ .fp4, .fp8, .fp8g, .bf16 }; +const modes = [_]Mode{ .fp4, .fp8, .mxfp8, .fp8g, .bf16 }; fn modeIndex(mode: Mode) usize { return std.mem.indexOfScalar(Mode, &modes, mode).?; @@ -29,10 +29,10 @@ fn modeIndex(mode: Mode) usize { pub fn sharedBytes(mode: Mode, bm: u32) u32 { const w: u32 = switch (mode) { .fp4 => bn * group / 2, - .fp8, .fp8g => bn * group, + .fp8, .mxfp8, .fp8g => bn * group, .bf16 => bn * group * 2, }; - const s: u32 = if (mode == .fp4 or mode == .fp8g) bn * 4 else 0; + const s: u32 = if (mode == .fp4 or mode == .fp8g) bn * 4 else if (mode == .mxfp8) bn * 2 else 0; const stage = (bm * group * 2 + w + s + 127) / 128 * 128; const partials = (bm / 16) * 2 * 4 * threads * 4; return @max(4 * stage, partials); @@ -186,6 +186,9 @@ test "tile shared bytes, slices and row tiles as the Python host picks them" { try std.testing.expectEqual(@as(u32, 25600), sharedBytes(.fp8g, 16)); try std.testing.expectEqual(@as(u32, 33792), sharedBytes(.fp8g, 32)); try std.testing.expectEqual(@as(u32, 50176), sharedBytes(.fp8g, 64)); + try std.testing.expectEqual(@as(u32, 25088), sharedBytes(.mxfp8, 16)); + try std.testing.expectEqual(@as(u32, 33280), sharedBytes(.mxfp8, 32)); + try std.testing.expectEqual(@as(u32, 49664), sharedBytes(.mxfp8, 64)); try std.testing.expectEqual(@as(u32, 17408), sharedBytes(.fp4, 16)); try std.testing.expectEqual(@as(u32, 25600), sharedBytes(.fp4, 32)); try std.testing.expectEqual(@as(u32, 41984), sharedBytes(.fp4, 64)); @@ -209,6 +212,7 @@ test "tile shared bytes, slices and row tiles as the Python host picks them" { } test "instance names are qmmf.cu's" { + try std.testing.expectEqualStrings("_ZN7tf_qmmf11qmmf_kernelILi2ELi16ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", symbol(.mxfp8, 16, true, false)); try std.testing.expectEqualStrings("_ZN7tf_qmmf11qmmf_kernelILi3ELi16ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", symbol(.fp8g, 16, true, false)); try std.testing.expectEqualStrings("_ZN7tf_qmmf11qmmf_kernelILi0ELi64ELi64ELi1ELi4ELi4ELb0ELb0ELb1EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", symbol(.fp4, 64, false, true)); } diff --git a/zig/src/cuda/root.zig b/zig/src/cuda/root.zig index 3053c09a6b..ada7a002ce 100644 --- a/zig/src/cuda/root.zig +++ b/zig/src/cuda/root.zig @@ -26,6 +26,7 @@ pub const aot = @import("aot.zig"); pub const kernels = @import("kernels.zig"); pub const qmmf = @import("qmmf.zig"); pub const fp8 = @import("fp8.zig"); +pub const mxfp8 = @import("mxfp8.zig"); pub const exlr8 = @import("exlr8.zig"); pub const nvfp4 = @import("nvfp4.zig"); pub const segments = @import("segments.zig"); @@ -51,6 +52,7 @@ test { _ = segments; _ = @import("qmmf.zig"); _ = @import("fp8.zig"); + _ = @import("mxfp8.zig"); _ = carveout; _ = @import("nvfp4.zig"); _ = @import("modelopt.zig"); diff --git a/zig/tests/cuda/main.zig b/zig/tests/cuda/main.zig index aec6b9ec1b..2d6b6bd85f 100644 --- a/zig/tests/cuda/main.zig +++ b/zig/tests/cuda/main.zig @@ -16,6 +16,7 @@ const prompt_cuts = @import("prompt_cuts.zig"); const chunk_costs = @import("chunk_costs.zig"); const grouped_tests = @import("grouped_tests.zig"); const fp8_tests = @import("fp8_tests.zig"); +const mxfp8_tests = @import("mxfp8_tests.zig"); const fp8_experts_tests = @import("fp8_experts_tests.zig"); const nvfp4_experts_tests = @import("nvfp4_experts_tests.zig"); const quant_ref_tests = @import("quant_ref_tests.zig"); @@ -61,6 +62,7 @@ const usage = \\ fp8-lane block-FP8 projections against the Python lane matmul's bytes (oracle/fp8_lane.py) \\ fp8-experts grouped block-FP8 experts against recorded fixture bytes \\ nvfp4-experts grouped NVFP4 experts against recorded fixture bytes + \\ mxfp8 MXFP8 strided projections against FP64 and row-count invariance \\ quant-ref FP8 and NVFP4 projections and experts against float64 references (synthetic) \\ carveout [MiB] [card] GB10 display memory: round trips and bandwidth (SKIP without the card) \\ nvfp4 NVFP4 projections against recorded lane-matmul and prompt-GEMM fixture bytes @@ -122,6 +124,7 @@ fn run(gpu: check.Gpu, cmd: []const u8, rest: []const [:0]const u8) !void { if (std.mem.eql(u8, cmd, "fp8-experts")) return fp8_experts_tests.experts(gpu, try arg(rest, 0)); if (std.mem.eql(u8, cmd, "nvfp4")) return nvfp4_tests.run(gpu, try arg(rest, 0)); if (std.mem.eql(u8, cmd, "nvfp4-experts")) return nvfp4_experts_tests.experts(gpu, try arg(rest, 0)); + if (std.mem.eql(u8, cmd, "mxfp8")) return mxfp8_tests.run(gpu); if (std.mem.eql(u8, cmd, "quant-ref")) return quant_ref_tests.run(gpu); if (std.mem.eql(u8, cmd, "modelopt")) return modelopt_tests.run(gpu); if (std.mem.eql(u8, cmd, "exlr8")) return exlr8_tests.run(gpu); diff --git a/zig/tests/cuda/mxfp8_tests.zig b/zig/tests/cuda/mxfp8_tests.zig new file mode 100644 index 0000000000..5417d3d1fd --- /dev/null +++ b/zig/tests/cuda/mxfp8_tests.zig @@ -0,0 +1,143 @@ +//! Native MXFP8: lossless packing, independent FP64 products, strided rows and row-count invariance. +const std = @import("std"); +const cuda = @import("cuda"); +const core = @import("core"); +const check = @import("check.zig"); + +fn bf(v: u16) f64 { + return @as(f32, @bitCast(@as(u32, v) << 16)); +} + +fn tobf(v: f32) u16 { + const u: u32 = @bitCast(v); + return @truncate((u + 0x7fff + ((u >> 16) & 1)) >> 16); +} + +pub fn run(gpu: check.Gpu) !void { + try scaleRange(gpu); + const a = gpu.gpa; + const n = 70; // a partial tile and padded rows + const k = 1024; // split-K, both cluster and explicit reduction + const m = 256; // fused prompt reduction + const ldx = k + 64; + const npad = cuda.mxfp8.padded(n); + var prng = std.Random.DefaultPrng.init(0x8e80_5eed); + const rng = prng.random(); + const codes = try a.alloc(u8, n * k); + defer a.free(codes); + for (codes) |*v| { + v.* = rng.int(u8); + if (v.* & 0x7f == 0x7f) v.* ^= 1; + } + const scales = try a.alloc(u8, n * (k / 32)); + defer a.free(scales); + for (scales) |*v| v.* = 121 + rng.uintLessThan(u8, 8); + // A dedicated row makes UE8M0 byte 0 observable in a normal BF16 result. + @memset(codes[0..k], 0x38); // E4M3 1 + @memset(scales[0 .. k / 32], 0); + const packed_codes = try a.alloc(u8, npad * k); + defer a.free(packed_codes); + cuda.mxfp8.packCodes(packed_codes, codes, n, k); + const packed_scales = try a.alloc(u8, npad * (k / 32)); + defer a.free(packed_scales); + cuda.mxfp8.packScales(packed_scales, scales, n, k); + var dw = try cuda.DeviceBuffer.fromHost(gpu.d, packed_codes); + defer dw.free(); + var ds = try cuda.DeviceBuffer.fromHost(gpu.d, packed_scales); + defer ds.free(); + const w = cuda.mxfp8.weight(dw.ptr, ds.ptr, n, k); + const x = try a.alloc(u16, m * ldx); + defer a.free(x); + @memset(x, 0x7fc0); // stride padding is poison, not another activation + for (0..m) |row| { + for (x[row * ldx ..][0..k]) |*v| v.* = tobf(rng.float(f32) * 2); + } + var dx = try cuda.DeviceBuffer.fromHost(gpu.d, std.mem.sliceAsBytes(x)); + defer dx.free(); + var dy = try cuda.DeviceBuffer.alloc(gpu.d, m * n * 2); + defer dy.free(); + var part = try cuda.DeviceBuffer.alloc(gpu.d, 8 * m * n * 4); + defer part.free(); + var lane = try cuda.qmmf.Lane.load(gpu.d, (try gpu.ctx.capability()) / 10); + defer lane.unload(); + var s = try cuda.Stream.init(gpu.d, false); + defer s.deinit(); + var baseline: ?[]u8 = null; + defer if (baseline) |b| a.free(b); + for ([_]bool{ true, false }) |clusters| { + lane.clusters = clusters and (try gpu.ctx.capability()) >= 90; + const lin: cuda.qlinear.Linear = .{ .lane = &lane }; + for ([_]usize{ 256, 1, 7, 17, 33 }) |rows| { + for ([_]bool{ false, true }) |prompt| { + try dy.fill8(0xff, s.handle); + const in: cuda.qlinear.Rows = .{ .x = dx.ptr, .ldx = ldx, .part = part.ptr }; + if (prompt) try lin.prompt(s, .{ .mxfp8 = w }, in, dy.ptr, rows) else try lin.decode(s, .{ .mxfp8 = w }, in, dy.ptr, rows); + try s.synchronize(); + const got = try a.alloc(u8, rows * n * 2); + defer a.free(got); + try dy.download(0, got); + if (baseline) |b| { + try check.sameBytes("MXFP8 row-count/dispatch/cluster invariance", got, b[0..got.len]); + } else baseline = try a.dupe(u8, got); + for (0..rows) |row| for (0..n) |col| { + var ref: f64 = 0; + var magnitude: f64 = 0; + for (0..k) |i| { + const scale = std.math.ldexp(@as(f64, 1), @as(i32, scales[col * (k / 32) + i / 32]) - 127); + const v = bf(x[row * ldx + i]) * core.quant.fp8.e4m3(codes[col * k + i]) * scale; + ref += v; + magnitude += @abs(v); + } + const actual = bf(std.mem.readInt(u16, got[(row * n + col) * 2 ..][0..2], .little)); + const tolerance = @abs(ref) / 128 + magnitude / 4096 + 0x1p-133; + try check.expect(@abs(actual - ref) <= tolerance, "MXFP8 FP64 row {d}, column {d}: {e} versus {e}, tolerance {e}", .{ row, col, actual, ref, tolerance }); + }; + } + } + } + check.pass("MXFP8: strided rows 1,7,17,33,256 equal across dispatch and reduction paths; FP64 including UE8M0 byte zero", .{}); +} + +/// Every UE8M0 encoding, through the production matmul: each dot product is exactly one times its scale. +fn scaleRange(gpu: check.Gpu) !void { + const a = gpu.gpa; + const n = 256; + const k = 64; + const codes = try a.alloc(u8, n * k); + defer a.free(codes); + @memset(codes, 0x38); + var scales: [n * 2]u8 = undefined; + for (0..n) |r| @memset(scales[r * 2 ..][0..2], @intCast(r)); + const tiled_codes = try a.alloc(u8, n * k); + defer a.free(tiled_codes); + cuda.mxfp8.packCodes(tiled_codes, codes, n, k); + var tiles: [n * 2]u8 = undefined; + cuda.mxfp8.packScales(&tiles, &scales, n, k); + var dw = try cuda.DeviceBuffer.fromHost(gpu.d, tiled_codes); + defer dw.free(); + var ds = try cuda.DeviceBuffer.fromHost(gpu.d, &tiles); + defer ds.free(); + var x: [k]u16 = @splat(0x3c80); // 1/64: 64 equal terms sum to one + var dx = try cuda.DeviceBuffer.fromHost(gpu.d, std.mem.sliceAsBytes(&x)); + defer dx.free(); + var dy = try cuda.DeviceBuffer.alloc(gpu.d, n * 2); + defer dy.free(); + var lane = try cuda.qmmf.Lane.load(gpu.d, (try gpu.ctx.capability()) / 10); + defer lane.unload(); + var s = try cuda.Stream.init(gpu.d, false); + defer s.deinit(); + try lane.matmul(s, dx.ptr, k, 1, cuda.mxfp8.weight(dw.ptr, ds.ptr, n, k), dy.ptr, 0); + try s.synchronize(); + var got: [n * 2]u8 = undefined; + try dy.download(0, &got); + for (0..n) |r| { + const bits = std.mem.readInt(u16, got[r * 2 ..][0..2], .little); + if (r == 255) { + try check.expect(std.math.isNan(bf(bits)), "UE8M0 reserved byte produces NaN", .{}); + } else { + const want: u16 = if (r == 0) 0x40 else @intCast(r << 7); + try check.expect(bits == want, "UE8M0 {d}: BF16 bits {x} versus {x}", .{ r, bits, want }); + } + } + check.pass("MXFP8: all 256 UE8M0 encodings preserve the specified value, including subnormal and NaN", .{}); +} From a008b5cd8157dba2827b254d3add4fe9a5f4b8e1 Mon Sep 17 00:00:00 2001 From: akol1 <15747200+akol1@users.noreply.github.com> Date: Fri, 16 Oct 2026 04:12:04 -0700 Subject: [PATCH 2/4] test(cuda): register MXFP8 launch resource limits --- zig/src/core/contract/registry.zig | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/zig/src/core/contract/registry.zig b/zig/src/core/contract/registry.zig index b79337b5c5..18d8928632 100644 --- a/zig/src/core/contract/registry.zig +++ b/zig/src/core/contract/registry.zig @@ -37,6 +37,14 @@ pub const cuda = [_]CudaRecord{ .{ .launch = .{ .name = "qmmf FP4 fp32 out 64-row tile", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 64, 0, 0 }), .image = "qmmf", .threads = 128, .dynamic = 41984 }, .usage = .{ .{ .regs = 148, .static = 0, .bound = 128 }, .{ .regs = 148, .static = 0, .bound = 128 }, .{ .regs = 180, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 } } }, .{ .launch = .{ .name = "qmmf FP4 fp32 out 64-row tile, cluster of 8", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 64, 1, 0 }), .image = "qmmf", .threads = 128, .dynamic = 41984, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 } } }, .{ .launch = .{ .name = "qmmf FP4 fp32 out fused slices", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 64, 0, 1 }), .image = "qmmf", .threads = 128, .dynamic = 41984 }, .usage = .{ .{ .regs = 182, .static = 0, .bound = 128 }, .{ .regs = 182, .static = 0, .bound = 128 }, .{ .regs = 228, .static = 0, .bound = 128 }, .{ .regs = 168, .static = 0, .bound = 128 }, .{ .regs = 168, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf MXFP8 16-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi16ELi64ELi1ELi4ELi4ELb0ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 25088 }, .usage = .{ .{ .regs = 94, .static = 0, .bound = 128 }, .{ .regs = 94, .static = 0, .bound = 128 }, .{ .regs = 94, .static = 0, .bound = 128 }, .{ .regs = 60, .static = 0, .bound = 128 }, .{ .regs = 60, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf MXFP8 16-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi16ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 25088, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 69, .static = 0, .bound = 128 }, .{ .regs = 69, .static = 0, .bound = 128 }, .{ .regs = 69, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf MXFP8 32-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi32ELi64ELi1ELi4ELi4ELb0ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 33280 }, .usage = .{ .{ .regs = 128, .static = 0, .bound = 128 }, .{ .regs = 128, .static = 0, .bound = 128 }, .{ .regs = 128, .static = 0, .bound = 128 }, .{ .regs = 57, .static = 0, .bound = 128 }, .{ .regs = 57, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf MXFP8 32-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi32ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 33280, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 66, .static = 0, .bound = 128 }, .{ .regs = 66, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf MXFP8 64-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi64ELi64ELi1ELi4ELi4ELb0ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49664, .opt_in = true }, .usage = .{ .{ .regs = 168, .static = 0, .bound = 128 }, .{ .regs = 168, .static = 0, .bound = 128 }, .{ .regs = 168, .static = 0, .bound = 128 }, .{ .regs = 108, .static = 0, .bound = 128 }, .{ .regs = 108, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf MXFP8 64-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi64ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49664, .opt_in = true, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 103, .static = 0, .bound = 128 }, .{ .regs = 103, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf MXFP8 64-row tile, fused slices", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi64ELi64ELi1ELi4ELi4ELb0ELb0ELb1EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49664, .opt_in = true }, .usage = .{ .{ .regs = 229, .static = 0, .bound = 128 }, .{ .regs = 229, .static = 0, .bound = 128 }, .{ .regs = 229, .static = 0, .bound = 128 }, .{ .regs = 158, .static = 0, .bound = 128 }, .{ .regs = 158, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf reduce_kernel", .symbol = "_ZN7tf_qmmf13reduce_kernelILb0EEEvPKfPvxif", .image = "qmmf", .threads = 256 }, .usage = .{ .{ .regs = 32, .static = 0, .bound = 0 }, .{ .regs = 36, .static = 0, .bound = 0 }, .{ .regs = 36, .static = 0, .bound = 0 }, .{ .regs = 40, .static = 0, .bound = 0 }, .{ .regs = 40, .static = 0, .bound = 0 } } }, .{ .launch = .{ .name = "qmmf reduce_kernel", .symbol = "_ZN7tf_qmmf13reduce_kernelILb1EEEvPKfPvxif", .image = "qmmf", .threads = 256 }, .usage = .{ .{ .regs = 32, .static = 0, .bound = 0 }, .{ .regs = 36, .static = 0, .bound = 0 }, .{ .regs = 36, .static = 0, .bound = 0 }, .{ .regs = 40, .static = 0, .bound = 0 }, .{ .regs = 40, .static = 0, .bound = 0 } } }, .{ .launch = .{ .name = "rank_sum tf_rank_sum", .symbol = "tf_rank_sum", .image = "rank_sum", .threads = 256 }, .usage = .{ .{ .regs = 32, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 } } }, .{ .launch = .{ .name = "vit vision_kernel<80, 8> (RADIO)", .symbol = vision80_symbol, .image = "prefill_attention", .threads = 256, .dynamic = 45056 }, .usage = .{ .{ .regs = 143, .static = 0, .bound = 256 }, .{ .regs = 143, .static = 0, .bound = 256 }, .{ .regs = 143, .static = 0, .bound = 256 }, .{ .regs = 122, .static = 0, .bound = 256 }, .{ .regs = 122, .static = 0, .bound = 256 } } }, From cb2c74984d4017b780f4e792001e1c3bffd99c59 Mon Sep 17 00:00:00 2001 From: akol1 <15747200+akol1@users.noreply.github.com> Date: Fri, 16 Oct 2026 06:07:35 -0700 Subject: [PATCH 3/4] cuda: refresh compiler-derived launch resource records --- docs/kernel-contract.md | 4 ++++ zig/src/core/contract/registry.zig | 12 ++++++------ 2 files changed, 10 insertions(+), 6 deletions(-) diff --git a/docs/kernel-contract.md b/docs/kernel-contract.md index e625bcb2fd..0b80446a42 100644 --- a/docs/kernel-contract.md +++ b/docs/kernel-contract.md @@ -108,3 +108,7 @@ CUDA: the RMS norm rows (16-byte form and the fp32-out tree), Nemotron's router tile, lane_gemv and its split-K form (both bit-equal to it), prompt16's FP4 tile with bf16 and fp32 out, qmmf's FP4 tiles with fp32 out (K slices met in a cluster, fused, or by `reduce_kernel`), and the reduce-scatter's `tf_rank_sum`. + +## CUDA compiler resource baseline + +The registered usage arrays were re-derived with NVCC 13.3.73 (`cuda_13.3.r13.3/compiler.38244171_0`) for SM 80, 86, 89, 120 and 121. `kernel-contract-limits` reads each compiled image and reports its register count, static shared memory and thread bound. Source changes or a different compiler may change these values; re-derive the arrays and run the live audit rather than changing the audit tolerance. Launch dimensions, dynamic shared memory, device limits and planted controls remain independent checks. Only SM 120 was executed on hardware in this update; the other four records come from their compiled images, not hardware execution. diff --git a/zig/src/core/contract/registry.zig b/zig/src/core/contract/registry.zig index b79337b5c5..028a30b511 100644 --- a/zig/src/core/contract/registry.zig +++ b/zig/src/core/contract/registry.zig @@ -20,20 +20,20 @@ const full72_symbol = "_ZN20tf_prefill_attention11full_kernelILi72ELi8ELi8EEEvPK /// CUDA kernels as launched; tf-cuda-test kernel-contract-limits re-derives the usage records. pub const cuda = [_]CudaRecord{ .{ .launch = .{ .name = "norms tf_rms_rows_bf16", .symbol = "tf_rms_rows_bf16", .image = "norms", .threads = 256 }, .usage = .{ .{ .regs = 79, .static = 32, .bound = 256 }, .{ .regs = 79, .static = 32, .bound = 256 }, .{ .regs = 79, .static = 32, .bound = 256 }, .{ .regs = 80, .static = 32, .bound = 256 }, .{ .regs = 80, .static = 32, .bound = 256 } } }, - .{ .launch = .{ .name = "norms tf_rms_rows", .symbol = "tf_rms_rows", .image = "norms", .threads = 1024 }, .usage = .{ .{ .regs = 32, .static = 132, .bound = 1024 }, .{ .regs = 46, .static = 132, .bound = 1024 }, .{ .regs = 46, .static = 132, .bound = 1024 }, .{ .regs = 45, .static = 132, .bound = 1024 }, .{ .regs = 45, .static = 132, .bound = 1024 } } }, + .{ .launch = .{ .name = "norms tf_rms_rows", .symbol = "tf_rms_rows", .image = "norms", .threads = 1024 }, .usage = .{ .{ .regs = 32, .static = 132, .bound = 1024 }, .{ .regs = 46, .static = 132, .bound = 1024 }, .{ .regs = 46, .static = 132, .bound = 1024 }, .{ .regs = 43, .static = 132, .bound = 1024 }, .{ .regs = 43, .static = 132, .bound = 1024 } } }, .{ .launch = .{ .name = "router tf_nemo_router", .symbol = "tf_nemo_router", .image = "nemotron_route", .threads = 256, .dynamic = 28800 }, .usage = .{ .{ .regs = 32, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 }, .{ .regs = 40, .static = 0, .bound = 256 } } }, - .{ .launch = .{ .name = "router tf_nemo_topk", .symbol = "tf_nemo_topk", .image = "nemotron_route", .threads = 32 }, .usage = .{ .{ .regs = 94, .static = 0, .bound = 32 }, .{ .regs = 95, .static = 0, .bound = 32 }, .{ .regs = 95, .static = 0, .bound = 32 }, .{ .regs = 91, .static = 0, .bound = 32 }, .{ .regs = 91, .static = 0, .bound = 32 } } }, + .{ .launch = .{ .name = "router tf_nemo_topk", .symbol = "tf_nemo_topk", .image = "nemotron_route", .threads = 32 }, .usage = .{ .{ .regs = 94, .static = 0, .bound = 32 }, .{ .regs = 95, .static = 0, .bound = 32 }, .{ .regs = 95, .static = 0, .bound = 32 }, .{ .regs = 96, .static = 0, .bound = 32 }, .{ .regs = 96, .static = 0, .bound = 32 } } }, .{ .launch = .{ .name = "qmm_group cluster of 8", .symbol = group_symbol, .image = "qmm_group", .threads = 128, .dynamic = 35328, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 } } }, .{ .launch = .{ .name = "qmm_group one CTA a tile", .symbol = group_symbol, .image = "qmm_group", .threads = 128, .dynamic = 35328 }, .usage = .{ .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 54, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 } } }, - .{ .launch = .{ .name = "lane_gemv", .symbol = gemv_symbol, .image = "lane_gemv", .threads = 128, .dynamic = 35328 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 57, .static = 0, .bound = 128 }, .{ .regs = 57, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 } } }, - .{ .launch = .{ .name = "lane_gemv split-K", .symbol = "tf_lane_gemv_split", .image = "lane_gemv", .threads = 128, .dynamic = 17664 }, .usage = .{ .{ .regs = 62, .static = 128, .bound = 128 }, .{ .regs = 56, .static = 128, .bound = 128 }, .{ .regs = 56, .static = 128, .bound = 128 }, .{ .regs = 44, .static = 128, .bound = 128 }, .{ .regs = 44, .static = 128, .bound = 128 } } }, + .{ .launch = .{ .name = "lane_gemv", .symbol = gemv_symbol, .image = "lane_gemv", .threads = 128, .dynamic = 35328 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 58, .static = 0, .bound = 128 }, .{ .regs = 58, .static = 0, .bound = 128 }, .{ .regs = 66, .static = 0, .bound = 128 }, .{ .regs = 66, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "lane_gemv split-K", .symbol = "tf_lane_gemv_split", .image = "lane_gemv", .threads = 128, .dynamic = 17664 }, .usage = .{ .{ .regs = 62, .static = 128, .bound = 128 }, .{ .regs = 56, .static = 128, .bound = 128 }, .{ .regs = 56, .static = 128, .bound = 128 }, .{ .regs = 72, .static = 128, .bound = 128 }, .{ .regs = 72, .static = 128, .bound = 128 } } }, .{ .launch = .{ .name = "prompt16 FP4", .symbol = std.fmt.comptimePrint(prompt_symbol, .{ 0, 0 }), .image = "prompt16", .threads = 128, .dynamic = 41984 }, .usage = .{ .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 } } }, - .{ .launch = .{ .name = "prompt16 FP8", .symbol = std.fmt.comptimePrint(prompt_symbol, .{ 1, 0 }), .image = "prompt16", .threads = 128, .dynamic = 49152 }, .usage = .{ .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 218, .static = 0, .bound = 128 }, .{ .regs = 218, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "prompt16 FP8", .symbol = std.fmt.comptimePrint(prompt_symbol, .{ 1, 0 }), .image = "prompt16", .threads = 128, .dynamic = 49152 }, .usage = .{ .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 243, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 } } }, .{ .launch = .{ .name = "prompt16 FP4 fp32 out", .symbol = std.fmt.comptimePrint(prompt_symbol, .{ 0, 1 }), .image = "prompt16", .threads = 128, .dynamic = 41984 }, .usage = .{ .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 255, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 }, .{ .regs = 212, .static = 0, .bound = 128 } } }, .{ .launch = .{ .name = "qmmf FP4 fp32 out 16-row tile", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 16, 0, 0 }), .image = "qmmf", .threads = 128, .dynamic = 17408 }, .usage = .{ .{ .regs = 88, .static = 0, .bound = 128 }, .{ .regs = 88, .static = 0, .bound = 128 }, .{ .regs = 96, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 } } }, .{ .launch = .{ .name = "qmmf FP4 fp32 out 16-row tile, cluster of 8", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 16, 1, 0 }), .image = "qmmf", .threads = 128, .dynamic = 17408, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 } } }, .{ .launch = .{ .name = "qmmf FP4 fp32 out 32-row tile", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 32, 0, 0 }), .image = "qmmf", .threads = 128, .dynamic = 25600 }, .usage = .{ .{ .regs = 106, .static = 0, .bound = 128 }, .{ .regs = 106, .static = 0, .bound = 128 }, .{ .regs = 128, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 } } }, - .{ .launch = .{ .name = "qmmf FP4 fp32 out 32-row tile, cluster of 8", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 32, 1, 0 }), .image = "qmmf", .threads = 128, .dynamic = 25600, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 63, .static = 0, .bound = 128 }, .{ .regs = 63, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP4 fp32 out 32-row tile, cluster of 8", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 32, 1, 0 }), .image = "qmmf", .threads = 128, .dynamic = 25600, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 61, .static = 0, .bound = 128 }, .{ .regs = 61, .static = 0, .bound = 128 } } }, .{ .launch = .{ .name = "qmmf FP4 fp32 out 64-row tile", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 64, 0, 0 }), .image = "qmmf", .threads = 128, .dynamic = 41984 }, .usage = .{ .{ .regs = 148, .static = 0, .bound = 128 }, .{ .regs = 148, .static = 0, .bound = 128 }, .{ .regs = 180, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 } } }, .{ .launch = .{ .name = "qmmf FP4 fp32 out 64-row tile, cluster of 8", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 64, 1, 0 }), .image = "qmmf", .threads = 128, .dynamic = 41984, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 } } }, .{ .launch = .{ .name = "qmmf FP4 fp32 out fused slices", .symbol = std.fmt.comptimePrint(qmmf_f32_symbol, .{ 64, 0, 1 }), .image = "qmmf", .threads = 128, .dynamic = 41984 }, .usage = .{ .{ .regs = 182, .static = 0, .bound = 128 }, .{ .regs = 182, .static = 0, .bound = 128 }, .{ .regs = 228, .static = 0, .bound = 128 }, .{ .regs = 168, .static = 0, .bound = 128 }, .{ .regs = 168, .static = 0, .bound = 128 } } }, From df44f518c49ab13e7154f7421f92d18af01667be Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Christian=20Schr=C3=B6der?= <1188270+cwschroeder@users.noreply.github.com> Date: Fri, 16 Oct 2026 13:19:46 +0200 Subject: [PATCH 4/4] Reject non-finite errors in CUDA reference comparisons --- zig/build/cuda.zig | 12 ++++ zig/tests/cuda/quant_ref_host_tests.zig | 83 +++++++++++++++++++++++++ zig/tests/cuda/quant_ref_tests.zig | 1 + 3 files changed, 96 insertions(+) create mode 100644 zig/tests/cuda/quant_ref_host_tests.zig diff --git a/zig/build/cuda.zig b/zig/build/cuda.zig index 684d6c3cc9..8aee541747 100644 --- a/zig/build/cuda.zig +++ b/zig/build/cuda.zig @@ -240,6 +240,18 @@ pub fn hostTests(b: *std.Build, draft_ids: *std.Build.Module, build_options: *st step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = b.createModule(.{ .root_source_file = b.path("zig/build/check_fatbins.zig"), .target = host }) })).step); const cuda = runtime(b, host, .debug, &.{}); const mods = family(b, host, .debug, cuda, draft_ids); + const reference_step = b.step("test-cuda-reference", "CUDA's FP64 comparison with CPU-only readback (no GPU)"); + for ([_]std.builtin.OptimizeMode{ .debug, .fast }) |mode| { + const reference = b.addRunArtifact(b.addTest(.{ .root_module = b.createModule(.{ + .root_source_file = b.path("zig/tests/cuda/quant_ref_host_tests.zig"), + .target = host, + .optimize = mode, + .link_libc = true, + .imports = &.{ .{ .name = "cuda", .module = cuda }, .{ .name = "core", .module = mods.core } }, + }) })); + step.dependOn(&reference.step); + reference_step.dependOn(&reference.step); + } const native_modules = engines(b, host, .debug, cuda, mods.core, mods.lanes, mods.nemotron, mods.heat); const native = native_modules.engines; // the HTTP server's unit tests (zig/src/server/root.zig), which the macOS build runs, on Linux hosts too diff --git a/zig/tests/cuda/quant_ref_host_tests.zig b/zig/tests/cuda/quant_ref_host_tests.zig new file mode 100644 index 0000000000..dc13f4235e --- /dev/null +++ b/zig/tests/cuda/quant_ref_host_tests.zig @@ -0,0 +1,83 @@ +//! CPU tests of Rig.close: only stream synchronization and readback are replaced; the comparison is the GPU gate's. +const std = @import("std"); +const cuda = @import("cuda"); +const Rig = @import("quant_ref_tests.zig").RigT; + +const Readback = struct { + var syncs: usize = 0; + var copies: usize = 0; + + fn sync(_: cuda.abi.Stream) callconv(.c) cuda.abi.Result { + syncs += 1; + return cuda.abi.success; + } + + fn copy(dst: ?*anyopaque, src: cuda.abi.DevicePtr, n: usize) callconv(.c) cuda.abi.Result { + const out: [*]u8 = @ptrCast(dst.?); + const in: [*]const u8 = @ptrFromInt(src); + @memcpy(out[0..n], in[0..n]); + copies += 1; + return cuda.abi.success; + } +}; + +fn close(comptime T: type, got: []const T, want: []const f64, tol: []const f64) !void { + var d: cuda.Driver = undefined; + d.api.cuStreamSynchronize = Readback.sync; + d.api.cuMemcpyDtoH_v2 = Readback.copy; + var arena = std.heap.ArenaAllocator.init(std.testing.allocator); + defer arena.deinit(); + var rig: Rig = .{ + .gpu = .{ .d = &d, .ctx = undefined, .gpa = std.testing.allocator, .io = std.testing.io }, + .a = arena.allocator(), + .s = .{ .d = &d, .handle = null }, + .rng = undefined, + }; + Readback.syncs = 0; + Readback.copies = 0; + const result = rig.close("CPU reference comparison", @intFromPtr(got.ptr), T == f32, want, tol); + try std.testing.expectEqual(@as(usize, 1), Readback.syncs); + try std.testing.expectEqual(@as(usize, 1), Readback.copies); + return result; +} + +const zero = [_]f64{ 0, 0, 0 }; +const tolerance = [_]f64{ 1, 1, 1 }; + +test "finite FP32 and BF16 outputs pass at the bound and fail beyond it" { + try close(f32, &.{ 0.5, 1, 0 }, &zero, &tolerance); + try close(u16, &.{ 0x3f00, 0x3f80, 0 }, &zero, &tolerance); + try std.testing.expectError(error.TestFailed, close(f32, &.{ 2, 0, 0 }, &zero, &tolerance)); + try std.testing.expectError(error.TestFailed, close(u16, &.{ 0x4000, 0, 0 }, &zero, &tolerance)); +} + +test "FP32 NaN cannot disappear before a later finite result" { + for ([_]u32{ 0x7fc00000, 0xffffffff }) |bits| { + for (0..3) |i| { + var got = [_]f32{ 0, 0, 0 }; + got[i] = @bitCast(bits); + try std.testing.expectError(error.TestFailed, close(f32, &got, &zero, &tolerance)); + } + } +} + +test "BF16 NaN cannot disappear before a later finite result" { + for ([_]u16{ 0x7fc0, 0xffff }) |bits| { + for (0..3) |i| { + var got = [_]u16{ 0, 0, 0 }; + got[i] = bits; + try std.testing.expectError(error.TestFailed, close(u16, &got, &zero, &tolerance)); + } + } +} + +test "NaN cannot erase an earlier out-of-bound or infinite result" { + try std.testing.expectError(error.TestFailed, close(f32, &.{ 0, std.math.inf(f32), 0 }, &zero, &tolerance)); + try std.testing.expectError(error.TestFailed, close(f32, &.{ 0, -std.math.inf(f32), 0 }, &zero, &tolerance)); + try std.testing.expectError(error.TestFailed, close(u16, &.{ 0, 0x7f80, 0 }, &zero, &tolerance)); + try std.testing.expectError(error.TestFailed, close(u16, &.{ 0, 0xff80, 0 }, &zero, &tolerance)); + try std.testing.expectError(error.TestFailed, close(f32, &.{ 2, std.math.nan(f32), 0 }, &zero, &tolerance)); + try std.testing.expectError(error.TestFailed, close(f32, &.{ std.math.inf(f32), std.math.nan(f32), 0 }, &zero, &tolerance)); + try std.testing.expectError(error.TestFailed, close(u16, &.{ 0x4000, 0xffff, 0 }, &zero, &tolerance)); + try std.testing.expectError(error.TestFailed, close(u16, &.{ 0x7f80, 0xffff, 0 }, &zero, &tolerance)); +} diff --git a/zig/tests/cuda/quant_ref_tests.zig b/zig/tests/cuda/quant_ref_tests.zig index 657cd388b7..a0f40b664c 100644 --- a/zig/tests/cuda/quant_ref_tests.zig +++ b/zig/tests/cuda/quant_ref_tests.zig @@ -58,6 +58,7 @@ const Rig = struct { for (want, tol, 0..) |w, t, i| { const g: f64 = if (f32_out) std.mem.bytesToValue(f32, raw[4 * i ..][0..4]) else bf(std.mem.bytesToValue(u16, raw[2 * i ..][0..2])); const q = @abs(g - w) / t; + try check.expect(std.math.isFinite(q), "{s}: non-finite error at element {d} (got {e}, reference {e}, tolerance {e})", .{ name, i, g, w, t }); if (!(q <= most)) { most = q; at = i;