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 01/10] 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 f6a0055f5d84dfe9f61342ef4923cd6c233bc037 Mon Sep 17 00:00:00 2001 From: akol1 <15747200+akol1@users.noreply.github.com> Date: Thu, 15 Oct 2026 21:23:49 -0700 Subject: [PATCH 02/10] Expose FP32 lane projection outputs before residual reduction --- zig/kernels/cuda/qmmf.cu | 36 +++++++ zig/src/cuda/qmmf.md | 19 ++++ zig/src/cuda/qmmf.zig | 40 ++++++-- zig/tests/cuda/main.zig | 3 + zig/tests/cuda/projection_f32_tests.zig | 127 ++++++++++++++++++++++++ 5 files changed, 214 insertions(+), 11 deletions(-) create mode 100644 zig/src/cuda/qmmf.md create mode 100644 zig/tests/cuda/projection_f32_tests.zig diff --git a/zig/kernels/cuda/qmmf.cu b/zig/kernels/cuda/qmmf.cu index 027b8ef630..f34fc02118 100644 --- a/zig/kernels/cuda/qmmf.cu +++ b/zig/kernels/cuda/qmmf.cu @@ -335,5 +335,41 @@ template __global__ void qmmf_kernel( 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); +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); +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); +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); +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); +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 reduce_kernel(const float* __restrict__, void* __restrict__, long long, int, float); template __global__ void reduce_kernel(const float* __restrict__, void* __restrict__, long long, int, float); } diff --git a/zig/src/cuda/qmmf.md b/zig/src/cuda/qmmf.md new file mode 100644 index 0000000000..82f5c5a1ad --- /dev/null +++ b/zig/src/cuda/qmmf.md @@ -0,0 +1,19 @@ +# FP32 projection outputs + +`cuda.qmmf.Lane.matmulF32` has the same inputs, packed weights, stride, +shape-driven split-K plan and scratch contract as `matmul`. Its output is +`[M, N]` FP32, requiring `M * N * 4` bytes. It preserves FP32 output for a +later residual or tensor-rank sum; the caller chooses when to round to BF16. + +All five lane formats are supported: NVFP4, tensor-scaled FP8, MXFP8, +block-scaled FP8 and BF16. The change exposes the existing kernel template's +FP32 output path. It does not change weight packing or accumulation order. +Rebuild the qmmf fatbin together with the native runtime: older fatbins do +not contain the FP32 instances or reduction symbol. + +`tf-cuda-test projection-f32` compares each format against an independent +FP64 calculation, checks conversion to the existing BF16 result, and varies +row count, activation stride, partial output tiles, fused reduction, +cluster reduction and explicit split reduction. The fixture is synthetic; +model correctness and multi-device collective behavior require separate +qualification. diff --git a/zig/src/cuda/qmmf.zig b/zig/src/cuda/qmmf.zig index 3e2b4168a7..adfddf7124 100644 --- a/zig/src/cuda/qmmf.zig +++ b/zig/src/cuda/qmmf.zig @@ -66,9 +66,9 @@ pub fn plan(m: usize, n: usize, k: usize, clusters: bool) Plan { } /// qmmf.cu's instance names, per mode: row tiles 16, 32, 64 (each plain, then with a cluster), then the fused kernel. -fn symbol(comptime mode: Mode, comptime bm: u32, comptime cluster: bool, comptime fused: bool) [:0]const u8 { - const name = "_ZN7tf_qmmf11qmmf_kernelILi{d}ELi{d}ELi64ELi1ELi4ELi4ELb0ELb{d}ELb{d}EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii"; - return std.fmt.comptimePrint(name, .{ @backingInt(mode), bm, @intFromBool(cluster), @intFromBool(fused) }); +fn symbol(comptime mode: Mode, comptime bm: u32, comptime cluster: bool, comptime fused: bool, comptime f32_out: bool) [:0]const u8 { + const name = "_ZN7tf_qmmf11qmmf_kernelILi{d}ELi{d}ELi64ELi1ELi4ELi4ELb{d}ELb{d}ELb{d}EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii"; + return std.fmt.comptimePrint(name, .{ @backingInt(mode), bm, @intFromBool(f32_out), @intFromBool(cluster), @intFromBool(fused) }); } pub const reduce_symbol = "_ZN7tf_qmmf13reduce_kernelILb0EEEvPKfPvxif"; @@ -78,22 +78,30 @@ pub const Lane = struct { mod: Module, fns: [modes.len][7]Function, // per mode: lane16, lane16c, lane32, lane32c, lane64, lane64c, fused reduce: Function, + f32_fns: [modes.len][7]Function, + reduce_f32: Function, clusters: bool, // sm_90 on: slices meet in a cluster's shared memory pub fn load(d: *const Driver, capability_major: u32) !Lane { if (!kernels.available) return error.BuiltWithoutKernels; var mod = try Module.load(d, kernels.qmmf); errdefer mod.unload(); - var l: Lane = .{ .mod = mod, .fns = undefined, .reduce = try mod.function(reduce_symbol), .clusters = capability_major >= 9 }; + var l: Lane = .{ .mod = mod, .fns = undefined, .f32_fns = undefined, .reduce = try mod.function(reduce_symbol), .reduce_f32 = try mod.function("_ZN7tf_qmmf13reduce_kernelILb1EEEvPKfPvxif"), .clusters = capability_major >= 9 }; inline for (modes, 0..) |mode, mi| { inline for (row_tiles, 0..) |bm, ti| inline for (.{ false, true }, 0..) |cl, ci| { - const f = try mod.function(symbol(mode, bm, cl, false)); + const f = try mod.function(symbol(mode, bm, cl, false, false)); try f.allowDynamicShared(sharedBytes(mode, bm)); l.fns[mi][ti * 2 + ci] = f; + const fp32_fn = try mod.function(symbol(mode, bm, cl, false, true)); + try fp32_fn.allowDynamicShared(sharedBytes(mode, bm)); + l.f32_fns[mi][ti * 2 + ci] = fp32_fn; }; - const f = try mod.function(symbol(mode, 64, false, true)); + const f = try mod.function(symbol(mode, 64, false, true, false)); try f.allowDynamicShared(sharedBytes(mode, 64)); l.fns[mi][6] = f; + const fp32_fn = try mod.function(symbol(mode, 64, false, true, true)); + try fp32_fn.allowDynamicShared(sharedBytes(mode, 64)); + l.f32_fns[mi][6] = fp32_fn; } return l; } @@ -110,8 +118,17 @@ pub const Lane = struct { /// y [m, n] bf16 = x [m, k] bf16 (rows `ldx` apart) times the weight, fp32 sums; `part` holds partBytes. pub fn matmul(l: Lane, s: Stream, x: u64, ldx: usize, m: usize, w: Weight, y: u64, part: u64) !void { + return l.matmulTo(s, x, ldx, m, w, y, part, false); + } + + /// The same accumulation and reduction, preserving FP32 outputs for a later residual or tensor-rank sum. + pub fn matmulF32(l: Lane, s: Stream, x: u64, ldx: usize, m: usize, w: Weight, y: u64, part: u64) !void { + return l.matmulTo(s, x, ldx, m, w, y, part, true); + } + + fn matmulTo(l: Lane, s: Stream, x: u64, ldx: usize, m: usize, w: Weight, y: u64, part: u64, f32_out: bool) !void { const p = plan(m, w.n, w.k, l.clusters); - const set = l.fns[modeIndex(w.mode)]; + const set = (if (f32_out) l.f32_fns else l.fns)[modeIndex(w.mode)]; const ti: usize = if (p.bm == 16) 0 else if (p.bm == 32) 1 else 2; const f = if (p.fused) set[6] else set[ti * 2 + @intFromBool(p.cluster)]; const rows_t = (m + p.bm - 1) / p.bm; @@ -144,7 +161,7 @@ pub const Lane = struct { r.add(@as(i64, @intCast(total))); r.add(@as(i32, @intCast(p.sk))); r.add(w.scale); - try launch_.launch(l.reduce, .{ .grid = .{ .x = @intCast((total + 255) / 256), .y = 1, .z = 1 }, .block = .{ .x = 256, .y = 1, .z = 1 } }, s, &r); + try launch_.launch(if (f32_out) l.reduce_f32 else l.reduce, .{ .grid = .{ .x = @intCast((total + 255) / 256), .y = 1, .z = 1 }, .block = .{ .x = 256, .y = 1, .z = 1 } }, s, &r); } } }; @@ -212,7 +229,8 @@ 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)); + try std.testing.expectEqualStrings("_ZN7tf_qmmf11qmmf_kernelILi4ELi64ELi64ELi1ELi4ELi4ELb1ELb0ELb1EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", symbol(.bf16, 64, false, true, true)); + try std.testing.expectEqualStrings("_ZN7tf_qmmf11qmmf_kernelILi2ELi16ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", symbol(.mxfp8, 16, true, false, false)); + try std.testing.expectEqualStrings("_ZN7tf_qmmf11qmmf_kernelILi3ELi16ELi64ELi1ELi4ELi4ELb0ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", symbol(.fp8g, 16, true, false, false)); + try std.testing.expectEqualStrings("_ZN7tf_qmmf11qmmf_kernelILi0ELi64ELi64ELi1ELi4ELi4ELb0ELb0ELb1EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", symbol(.fp4, 64, false, true, false)); } diff --git a/zig/tests/cuda/main.zig b/zig/tests/cuda/main.zig index 2d6b6bd85f..ac5249e1d6 100644 --- a/zig/tests/cuda/main.zig +++ b/zig/tests/cuda/main.zig @@ -17,6 +17,7 @@ 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 projection_f32_tests = @import("projection_f32_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"); @@ -62,6 +63,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 + \\ projection-f32 FP32 projection outputs against FP64, rounding and row/reduction invariance \\ 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) @@ -124,6 +126,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, "projection-f32")) return projection_f32_tests.run(gpu); 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); diff --git a/zig/tests/cuda/projection_f32_tests.zig b/zig/tests/cuda/projection_f32_tests.zig new file mode 100644 index 0000000000..85b1160365 --- /dev/null +++ b/zig/tests/cuda/projection_f32_tests.zig @@ -0,0 +1,127 @@ +//! FP32 lane output before residual/rank combination: independent FP64, rounding and row/reduction invariance. +const std = @import("std"); +const cuda = @import("cuda"); +const core = @import("core"); +const check = @import("check.zig"); +const reference = @import("quant_ref_tests.zig"); +const Rig = reference.RigT; +const bf = reference.helpers.bfOf; +const bf16 = reference.helpers.bf16Of; +const Mode = cuda.qmmf.Mode; + +pub fn run(gpu: check.Gpu) !void { + var arena = std.heap.ArenaAllocator.init(gpu.gpa); + defer arena.deinit(); + var s = try cuda.Stream.init(gpu.d, false); + defer s.deinit(); + var prng = std.Random.DefaultPrng.init(0xf32_5eed); + var r: Rig = .{ .gpu = gpu, .a = arena.allocator(), .s = s, .rng = prng.random() }; + defer { + for (r.bufs.items) |*b| b.free(); + r.bufs.deinit(gpu.gpa); + } + var lane = try cuda.qmmf.Lane.load(gpu.d, (try gpu.ctx.capability()) / 10); + defer lane.unload(); + for ([_]Mode{ .fp4, .fp8, .mxfp8, .fp8g, .bf16 }) |mode| try projection(&r, &lane, mode); +} + +fn projection(r: *Rig, lane: *cuda.qmmf.Lane, mode: Mode) !void { + const n = 70; + const k = 1024; + const m = 256; + const ldx = k + 64; + const npad = cuda.fp8.padded(n); + const a = r.a; + const codes = try a.alloc(u8, n * k); + for (codes) |*v| v.* = reference.helpers.fp8CodeOf(r); + const scales = try a.alloc(u8, n * (k / 16)); + for (scales) |*v| v.* = if (mode == .mxfp8) 121 + r.rng.uintLessThan(u8, 8) else 0x20 + r.rng.uintLessThan(u8, 0x28); + const inv = try a.alloc(f32, ((n + 127) / 128) * (k / 128)); + for (inv) |*v| v.* = r.rng.float(f32) * 0.01 + 0.001; + const w16 = try r.bfs(n * k, 0.25); + const packed16 = try a.alloc(u16, npad * k); + const tiled_codes = std.mem.sliceAsBytes(packed16); + const global: f32 = 0.0123; + var w: cuda.qmmf.Weight = undefined; + switch (mode) { + .bf16 => { + cuda.qmmf.packBf16(packed16, w16, n, k); + w = cuda.modelopt.bf16Weight(try r.dev(tiled_codes), n, k); + }, + .fp4 => { + const words = try a.alloc(u32, npad * k / 8); + cuda.nvfp4.packWords(words, codes[0 .. n * k / 2], n, k); + const tiles = try a.alloc(u8, npad * (k / 16)); + cuda.nvfp4.packScales(tiles, scales, n, k); + w = cuda.nvfp4.weight(try r.dev(std.mem.sliceAsBytes(words)), try r.dev(tiles), global, n, k); + }, + .fp8, .fp8g, .mxfp8 => { + cuda.fp8.packCodes(tiled_codes[0 .. npad * k], codes, n, k); + const dw = try r.dev(tiled_codes[0 .. npad * k]); + w = switch (mode) { + .fp8 => cuda.modelopt.fp8Weight(dw, global, n, k), + .fp8g => blk: { + const tiles = try a.alloc(f32, npad * (k / 64)); + cuda.fp8.packScales(tiles, inv, n, k); + break :blk cuda.fp8.weight(dw, try r.dev(std.mem.sliceAsBytes(tiles)), n, k); + }, + .mxfp8 => blk: { + const tiles = try a.alloc(u8, npad * (k / 32)); + cuda.mxfp8.packScales(tiles, scales[0 .. n * k / 32], n, k); + break :blk cuda.mxfp8.weight(dw, try r.dev(tiles), n, k); + }, + else => unreachable, + }; + }, + } + const x = try a.alloc(u16, m * ldx); + @memset(x, 0x7fc0); + for (0..m) |row| { + for (x[row * ldx ..][0..k]) |*v| v.* = bf16((r.rng.float(f32) - 0.5) * 2); + } + const dx = try r.dev(std.mem.sliceAsBytes(x)); + const out = try r.zeros(m * n * 4); + const out_bf16 = try r.zeros(m * n * 2); + const part = try r.zeros(8 * m * n * 4); + const want = try a.alloc(f64, m * n); + const tol = try a.alloc(f64, m * n); + for (0..m) |row| for (0..n) |col| { + var sum: f64 = 0; + var magnitude: f64 = 0; + for (0..k) |i| { + const code = codes[col * k + i]; + const weight: f64 = switch (mode) { + .bf16 => bf(w16[col * k + i]), + .fp8 => core.quant.fp8.e4m3(code) * global, + .fp8g => core.quant.fp8.e4m3(code) * inv[(col / 128) * (k / 128) + i / 128], + .mxfp8 => core.quant.fp8.e4m3(code) * std.math.ldexp(@as(f64, 1), @as(i32, scales[col * (k / 32) + i / 32]) - 127), + .fp4 => blk: { + const byte = codes[col * (k / 2) + i / 2]; + break :blk core.quant.nvfp4.e2m1(if (i % 2 == 0) byte & 15 else byte >> 4) * core.quant.fp8.e4m3(scales[col * (k / 16) + i / 16]) * global; + }, + }; + const v = bf(x[row * ldx + i]) * weight; + sum += v; + magnitude += @abs(v); + } + want[row * n + col] = sum; + tol[row * n + col] = magnitude * 2e-6 + @abs(sum) * 1e-6 + 1e-30; + }; + var baseline: ?[]u8 = null; + for ([_]bool{ true, false }) |clusters| { + lane.clusters = clusters and (try r.gpu.ctx.capability()) >= 90; + for ([_]usize{ 256, 1, 7, 17, 33 }) |rows| { + try lane.matmulF32(r.s, dx, ldx, rows, w, out, part); + try r.close("FP32 projection", out, true, want[0 .. rows * n], tol[0 .. rows * n]); + const got = try r.back(out, rows * n * 4); + if (baseline) |b| try check.sameBytes("FP32 row/reduction invariance", got, b[0..got.len]) else baseline = got; + try lane.matmul(r.s, dx, ldx, rows, w, out_bf16, part); + const rounded = try r.back(out_bf16, rows * n * 2); + for (0..rows * n) |i| { + const v = std.mem.readInt(u32, got[i * 4 ..][0..4], .little); + try check.expect(bf16(@bitCast(v)) == std.mem.readInt(u16, rounded[i * 2 ..][0..2], .little), "FP32 rounds once to the existing BF16 result", .{}); + } + } + } + check.pass("{t}: FP32 outputs preserve FP64 accuracy, BF16 rounding and row/reduction invariance", .{mode}); +} From 50c4f56ff5622829d3067e5bd77a173fa139611f Mon Sep 17 00:00:00 2001 From: akol1 <15747200+akol1@users.noreply.github.com> Date: Fri, 16 Oct 2026 03:13:42 -0700 Subject: [PATCH 03/10] Exercise FP32 packed projections through numerical contract gates --- zig/tests/cuda/contract_projection_f32.zig | 157 +++++++++++++++++++++ zig/tests/cuda/contract_tests.zig | 8 ++ 2 files changed, 165 insertions(+) create mode 100644 zig/tests/cuda/contract_projection_f32.zig diff --git a/zig/tests/cuda/contract_projection_f32.zig b/zig/tests/cuda/contract_projection_f32.zig new file mode 100644 index 0000000000..26d32ac369 --- /dev/null +++ b/zig/tests/cuda/contract_projection_f32.zig @@ -0,0 +1,157 @@ +//! Packed projections writing FP32, with independent FP64 products and mirrored accumulator stress. +const std = @import("std"); +const cuda = @import("cuda"); +const core = @import("core"); +const contract = core.contract; +const c = contract.spec; +const g = contract.gates; +const Dense = contract.dense.Dense; +const Rig = @import("contract_rig.zig").Rig; + +pub const modes = [_]cuda.qmmf.Mode{ .bf16, .fp8, .fp8g, .mxfp8 }; +const served = [_]c.Shape{ + .{ .rows = 1, .n = 70, .k = 1024 }, + .{ .rows = 17, .n = 70, .k = 1024 }, + .{ .rows = 33, .n = 70, .k = 1024 }, + .{ .rows = 256, .n = 70, .k = 1024 }, + .{ .rows = 16, .n = 1280, .k = 2560 }, + .{ .rows = 16, .n = 2560, .k = 640 }, +}; + +fn depth(shape: c.Shape) usize { + return shape.k / 16 + 24; +} + +pub const Product = struct { + rig: Rig, + lane: *const cuda.qmmf.Lane, + mode: cuda.qmmf.Mode, + clusters: bool, + d: ?Dense = null, + weight: cuda.qmmf.Weight = undefined, + x: u64 = 0, + out: u64 = 0, + part: u64 = 0, + + pub fn deinit(self: *Product) void { + if (self.d) |d| d.deinit(); + self.rig.deinit(); + } + + pub fn kernel(self: *Product) g.Kernel { + const name = switch (self.mode) { + .bf16 => if (self.clusters) "cuda BF16 fp32 clustered" else "cuda BF16 fp32 reduce", + .fp8 => if (self.clusters) "cuda FP8 fp32 clustered" else "cuda FP8 fp32 reduce", + .fp8g => if (self.clusters) "cuda block FP8 fp32 clustered" else "cuda block FP8 fp32 reduce", + .mxfp8 => if (self.clusters) "cuda MXFP8 fp32 clustered" else "cuda MXFP8 fp32 reduce", + else => unreachable, + }; + return .{ .spec = .{ + .name = name, + .accum = .f32, + .depth = depth, + .out = .f32, + .tiles = .{ .{ .size = 16 }, .{ .size = 64 }, .{ .size = 512, .step = 128, .min = 128 } }, + .served = &served, + }, .ctx = self, .vt = &vt }; + } + + const vt = g.hooks(Product); + + pub fn prepare(self: *Product, shape: c.Shape, fill: g.Fill, seed: u64) !bool { + if (self.d) |d| d.deinit(); + self.d = null; + self.rig.reset(); + self.part = 0; + if (shape.rows == 0 or shape.n == 0 or shape.k == 0 or shape.k % 128 != 0) return false; + const d = try Dense.init(self.rig.gpu.gpa, shape, 1); + self.d = d; + d.fillX(fill, seed, 16); + try self.prepareWeight(shape, fill, seed); + self.x = try self.rig.input(std.mem.sliceAsBytes(d.x)); + const bytes = self.lane.partBytes(shape.rows, self.weight); + if (bytes > 0) self.part = try self.rig.scratch(bytes); + self.out = try self.rig.output(shape.rows * shape.n * 4); + return true; + } + + fn prepareWeight(self: *Product, shape: c.Shape, fill: g.Fill, seed: u64) !void { + var arena = std.heap.ArenaAllocator.init(self.rig.gpu.gpa); + defer arena.deinit(); + const a = arena.allocator(); + const d = self.d.?; + const n = shape.n; + const k = shape.k; + const npad = cuda.fp8.padded(n); + var prng = std.Random.DefaultPrng.init(seed ^ 0xf32); + const random = prng.random(); + const codes = try a.alloc(u8, n * k); + const words = try a.alloc(u16, n * k); + const ue8 = try a.alloc(u8, n * k / 32); + for (ue8) |*v| v.* = if (fill == .stress) 127 else 123 + random.uintLessThan(u8, 8); + const inv = try a.alloc(f32, ((n + 127) / 128) * (k / 128)); + for (inv) |*v| v.* = if (fill == .stress) 1 else std.math.ldexp(@as(f32, 1), @as(i32, random.uintLessThan(u8, 8)) - 4); + for (0..n) |j| for (0..k) |i| { + const at = j * k + i; + codes[at] = if (fill == .stress) 0x38 + random.uintLessThan(u8, 24) else random.uintLessThan(u8, 0x7f); + if (fill != .stress and random.boolean()) codes[at] |= 0x80; + words[at] = c.toBf16(if (fill == .stress) @floatFromInt(1 + random.uintLessThan(u8, 7)) else random.float(f32) * 4 - 2); + }; + if (fill == .stress) for (0..n) |j| for (k / 2..k) |i| { + codes[j * k + i] = codes[j * k + i - k / 2] | 0x80; + words[j * k + i] = words[j * k + i - k / 2] | 0x8000; + }; + const scale: f32 = if (fill == .stress) 1 else 0.125; + for (0..n) |j| for (0..k) |i| { + const value = core.quant.fp8.e4m3(codes[j * k + i]); + d.w[j * k + i] = switch (self.mode) { + .bf16 => c.fromBf16(words[j * k + i]), + .fp8 => value * scale, + .fp8g => value * inv[(j / 128) * (k / 128) + i / 128], + .mxfp8 => value * std.math.ldexp(@as(f64, 1), @as(i32, ue8[j * (k / 32) + i / 32]) - 127), + else => unreachable, + }; + }; + if (self.mode == .bf16) { + const tiles = try a.alloc(u16, npad * k); + cuda.qmmf.packBf16(tiles, words, n, k); + self.weight = cuda.modelopt.bf16Weight(try self.rig.input(std.mem.sliceAsBytes(tiles)), n, k); + return; + } + const tiles = try a.alloc(u8, npad * k); + cuda.fp8.packCodes(tiles, codes, n, k); + const dw = try self.rig.input(tiles); + self.weight = switch (self.mode) { + .fp8 => cuda.modelopt.fp8Weight(dw, scale, n, k), + .mxfp8 => blk: { + const scales = try a.alloc(u8, npad * k / 32); + cuda.mxfp8.packScales(scales, ue8, n, k); + break :blk cuda.mxfp8.weight(dw, try self.rig.input(scales), n, k); + }, + .fp8g => blk: { + const scales = try a.alloc(f32, npad * k / 64); + cuda.fp8.packScales(scales, inv, n, k); + break :blk cuda.fp8.weight(dw, try self.rig.input(std.mem.sliceAsBytes(scales)), n, k); + }, + else => unreachable, + }; + } + + pub fn run(self: *Product) !void { + const shape = self.d.?.shape; + try self.lane.matmulOut(self.rig.s, self.x, shape.k, shape.rows, self.weight, self.out, .f32, self.part); + try self.rig.s.synchronize(); + } + + pub fn count(self: *Product) usize { + return self.d.?.outputs(); + } + + pub fn reference(self: *Product, want: []f64, magnitude: []f64) !void { + try self.d.?.reference(want, magnitude); + } + + pub fn simulate(self: *Product, accumulator: c.Accum, order: g.Order, out: []f64) !void { + try self.d.?.simulate(accumulator, order, out); + } +}; diff --git a/zig/tests/cuda/contract_tests.zig b/zig/tests/cuda/contract_tests.zig index c89c814c6c..e1f58336a2 100644 --- a/zig/tests/cuda/contract_tests.zig +++ b/zig/tests/cuda/contract_tests.zig @@ -8,6 +8,7 @@ const rows_k = @import("contract_rows.zig"); const router_k = @import("contract_router.zig"); const mm_k = @import("contract_mm.zig"); const mx_k = @import("contract_mxfp8.zig"); +const wide_k = @import("contract_projection_f32.zig"); const rank_k = @import("contract_rank.zig"); const contract = core.contract; @@ -93,6 +94,13 @@ pub fn run(gpu: check.Gpu) !void { }; defer for (&mx) |*m| m.deinit(); for (&mx) |*m| try kernels.append(a, m.kernel()); + var wide: [wide_k.modes.len * 2]wide_k.Product = undefined; + for (wide_k.modes, 0..) |mode, i| { + wide[i * 2] = .{ .rig = .{ .gpu = gpu, .s = s }, .lane = &lane, .mode = mode, .clusters = true }; + wide[i * 2 + 1] = .{ .rig = .{ .gpu = gpu, .s = s }, .lane = &lane_reduce, .mode = mode, .clusters = false }; + } + defer for (&wide) |*m| m.deinit(); + for (&wide) |*m| try kernels.append(a, m.kernel()); var rank = rank_k.Rank{ .rig = .{ .gpu = gpu, .s = s }, .sum = &rank_sum }; defer rank.deinit(); try kernels.append(a, rank.kernel()); From 37578da9e2f739c6981999b6db0b4a61f572b206 Mon Sep 17 00:00:00 2001 From: akol1 <15747200+akol1@users.noreply.github.com> Date: Fri, 16 Oct 2026 03:16:59 -0700 Subject: [PATCH 04/10] Cover served dense and attention projection dimensions --- zig/tests/cuda/contract_projection_f32.zig | 2 ++ 1 file changed, 2 insertions(+) diff --git a/zig/tests/cuda/contract_projection_f32.zig b/zig/tests/cuda/contract_projection_f32.zig index 26d32ac369..1e172dec30 100644 --- a/zig/tests/cuda/contract_projection_f32.zig +++ b/zig/tests/cuda/contract_projection_f32.zig @@ -16,6 +16,8 @@ const served = [_]c.Shape{ .{ .rows = 256, .n = 70, .k = 1024 }, .{ .rows = 16, .n = 1280, .k = 2560 }, .{ .rows = 16, .n = 2560, .k = 640 }, + .{ .rows = 7, .n = 2560, .k = 10240 }, + .{ .rows = 7, .n = 13952, .k = 2560 }, }; fn depth(shape: c.Shape) usize { 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 05/10] 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 5bb739249784564c9b8ec591c94b8c4aa8fbd5ad Mon Sep 17 00:00:00 2001 From: akol1 <15747200+akol1@users.noreply.github.com> Date: Fri, 16 Oct 2026 04:42:24 -0700 Subject: [PATCH 06/10] test(cuda): register FP32 projection launch resources --- zig/src/core/contract/registry.zig | 28 ++++++++++++++++++++++++++++ 1 file changed, 28 insertions(+) diff --git a/zig/src/core/contract/registry.zig b/zig/src/core/contract/registry.zig index 18d8928632..a14c3f58a8 100644 --- a/zig/src/core/contract/registry.zig +++ b/zig/src/core/contract/registry.zig @@ -45,6 +45,34 @@ pub const cuda = [_]CudaRecord{ .{ .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 FP8 fp32 out 16-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi1ELi16ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 24576 }, .usage = .{ .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 50, .static = 0, .bound = 128 }, .{ .regs = 50, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8 fp32 out 16-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi1ELi16ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 24576, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 40, .static = 0, .bound = 128 }, .{ .regs = 40, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8 fp32 out 32-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi1ELi32ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 32768 }, .usage = .{ .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8 fp32 out 32-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi1ELi32ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 32768, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8 fp32 out 64-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi1ELi64ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49152 }, .usage = .{ .{ .regs = 126, .static = 0, .bound = 128 }, .{ .regs = 126, .static = 0, .bound = 128 }, .{ .regs = 126, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8 fp32 out 64-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi1ELi64ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49152, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8 fp32 out 64-row tile, fused slices", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi1ELi64ELi64ELi1ELi4ELi4ELb1ELb0ELb1EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49152 }, .usage = .{ .{ .regs = 179, .static = 0, .bound = 128 }, .{ .regs = 179, .static = 0, .bound = 128 }, .{ .regs = 179, .static = 0, .bound = 128 }, .{ .regs = 140, .static = 0, .bound = 128 }, .{ .regs = 140, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf MXFP8 fp32 out 16-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi16ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 25088 }, .usage = .{ .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 60, .static = 0, .bound = 128 }, .{ .regs = 60, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf MXFP8 fp32 out 16-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi16ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__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 fp32 out 32-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi32ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__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 fp32 out 32-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi32ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 33280, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 66, .static = 0, .bound = 128 }, .{ .regs = 66, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf MXFP8 fp32 out 64-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi64ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__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 fp32 out 64-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi64ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__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 fp32 out 64-row tile, fused slices", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi2ELi64ELi64ELi1ELi4ELi4ELb1ELb0ELb1EEEvPK13__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 FP8G fp32 out 16-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi3ELi16ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 25600 }, .usage = .{ .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 92, .static = 0, .bound = 128 }, .{ .regs = 58, .static = 0, .bound = 128 }, .{ .regs = 58, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8G fp32 out 16-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi3ELi16ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .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 = 64, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8G fp32 out 32-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi3ELi32ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 33792 }, .usage = .{ .{ .regs = 116, .static = 0, .bound = 128 }, .{ .regs = 116, .static = 0, .bound = 128 }, .{ .regs = 116, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 }, .{ .regs = 56, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8G fp32 out 32-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi3ELi32ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 33792, .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 = 57, .static = 0, .bound = 128 }, .{ .regs = 57, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8G fp32 out 64-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi3ELi64ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 50176, .opt_in = true }, .usage = .{ .{ .regs = 164, .static = 0, .bound = 128 }, .{ .regs = 164, .static = 0, .bound = 128 }, .{ .regs = 164, .static = 0, .bound = 128 }, .{ .regs = 108, .static = 0, .bound = 128 }, .{ .regs = 108, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8G fp32 out 64-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi3ELi64ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 50176, .opt_in = true, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 78, .static = 0, .bound = 128 }, .{ .regs = 122, .static = 0, .bound = 128 }, .{ .regs = 122, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf FP8G fp32 out 64-row tile, fused slices", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi3ELi64ELi64ELi1ELi4ELi4ELb1ELb0ELb1EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 50176, .opt_in = true }, .usage = .{ .{ .regs = 196, .static = 0, .bound = 128 }, .{ .regs = 196, .static = 0, .bound = 128 }, .{ .regs = 196, .static = 0, .bound = 128 }, .{ .regs = 155, .static = 0, .bound = 128 }, .{ .regs = 155, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf BF16 fp32 out 16-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi4ELi16ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 40960 }, .usage = .{ .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 80, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf BF16 fp32 out 16-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi4ELi16ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 40960, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 40, .static = 0, .bound = 128 }, .{ .regs = 40, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf BF16 fp32 out 32-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi4ELi32ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49152 }, .usage = .{ .{ .regs = 93, .static = 0, .bound = 128 }, .{ .regs = 93, .static = 0, .bound = 128 }, .{ .regs = 93, .static = 0, .bound = 128 }, .{ .regs = 55, .static = 0, .bound = 128 }, .{ .regs = 55, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf BF16 fp32 out 32-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi4ELi32ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 49152, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 }, .{ .regs = 48, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf BF16 fp32 out 64-row tile", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi4ELi64ELi64ELi1ELi4ELi4ELb1ELb0ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 65536, .opt_in = true }, .usage = .{ .{ .regs = 124, .static = 0, .bound = 128 }, .{ .regs = 124, .static = 0, .bound = 128 }, .{ .regs = 124, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 }, .{ .regs = 72, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf BF16 fp32 out 64-row tile, cluster of 8", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi4ELi64ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 65536, .opt_in = true, .cluster = 8, .min_sm = 90 }, .usage = .{ .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 }, .{ .regs = 64, .static = 0, .bound = 128 } } }, + .{ .launch = .{ .name = "qmmf BF16 fp32 out 64-row tile, fused slices", .symbol = "_ZN7tf_qmmf11qmmf_kernelILi4ELi64ELi64ELi1ELi4ELi4ELb1ELb0ELb1EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", .image = "qmmf", .threads = 128, .dynamic = 65536, .opt_in = true }, .usage = .{ .{ .regs = 180, .static = 0, .bound = 128 }, .{ .regs = 180, .static = 0, .bound = 128 }, .{ .regs = 180, .static = 0, .bound = 128 }, .{ .regs = 140, .static = 0, .bound = 128 }, .{ .regs = 140, .static = 0, .bound = 128 } } }, .{ .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 07/10] 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 08/10] 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; From 37359a17e1be44fd50cd0d7b4086b07a03314601 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 12:41:51 +0200 Subject: [PATCH 09/10] test: poison each FP32 projection dispatch and reduction scratch --- zig/tests/cuda/projection_f32_tests.zig | 16 +++++++++++++--- 1 file changed, 13 insertions(+), 3 deletions(-) diff --git a/zig/tests/cuda/projection_f32_tests.zig b/zig/tests/cuda/projection_f32_tests.zig index 85b1160365..19a62fdf4f 100644 --- a/zig/tests/cuda/projection_f32_tests.zig +++ b/zig/tests/cuda/projection_f32_tests.zig @@ -80,9 +80,15 @@ fn projection(r: *Rig, lane: *cuda.qmmf.Lane, mode: Mode) !void { for (x[row * ldx ..][0..k]) |*v| v.* = bf16((r.rng.float(f32) - 0.5) * 2); } const dx = try r.dev(std.mem.sliceAsBytes(x)); - const out = try r.zeros(m * n * 4); - const out_bf16 = try r.zeros(m * n * 2); - const part = try r.zeros(8 * m * n * 4); + var out_buf = try cuda.DeviceBuffer.alloc(r.gpu.d, m * n * 4); + defer out_buf.free(); + var bf16_buf = try cuda.DeviceBuffer.alloc(r.gpu.d, m * n * 2); + defer bf16_buf.free(); + var part_buf = try cuda.DeviceBuffer.alloc(r.gpu.d, 8 * m * n * 4); + defer part_buf.free(); + const out = out_buf.ptr; + const out_bf16 = bf16_buf.ptr; + const part = part_buf.ptr; const want = try a.alloc(f64, m * n); const tol = try a.alloc(f64, m * n); for (0..m) |row| for (0..n) |col| { @@ -111,10 +117,14 @@ fn projection(r: *Rig, lane: *cuda.qmmf.Lane, mode: Mode) !void { for ([_]bool{ true, false }) |clusters| { lane.clusters = clusters and (try r.gpu.ctx.capability()) >= 90; for ([_]usize{ 256, 1, 7, 17, 33 }) |rows| { + try out_buf.fill8(0xff, r.s.handle); + try part_buf.fill8(0xff, r.s.handle); try lane.matmulF32(r.s, dx, ldx, rows, w, out, part); try r.close("FP32 projection", out, true, want[0 .. rows * n], tol[0 .. rows * n]); const got = try r.back(out, rows * n * 4); if (baseline) |b| try check.sameBytes("FP32 row/reduction invariance", got, b[0..got.len]) else baseline = got; + try bf16_buf.fill8(0xff, r.s.handle); + try part_buf.fill8(0xff, r.s.handle); try lane.matmul(r.s, dx, ldx, rows, w, out_bf16, part); const rounded = try r.back(out_bf16, rows * n * 2); for (0..rows * n) |i| { From 4544c5c5da7c5a9b6ffa0cfa46275e32cb011397 Mon Sep 17 00:00:00 2001 From: akol1 <15747200+akol1@users.noreply.github.com> Date: Fri, 16 Oct 2026 14:08:15 -0700 Subject: [PATCH 10/10] test(cuda): reject no-op projection dispatches after stale prefixes --- zig/build/cuda.zig | 10 ++ zig/tests/cuda/projection_f32_host_tests.zig | 109 +++++++++++++++++++ zig/tests/cuda/projection_f32_tests.zig | 16 +-- 3 files changed, 128 insertions(+), 7 deletions(-) create mode 100644 zig/tests/cuda/projection_f32_host_tests.zig diff --git a/zig/build/cuda.zig b/zig/build/cuda.zig index 8aee541747..ee6685bf9d 100644 --- a/zig/build/cuda.zig +++ b/zig/build/cuda.zig @@ -241,6 +241,7 @@ pub fn hostTests(b: *std.Build, draft_ids: *std.Build.Module, build_options: *st 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)"); + const projection_step = b.step("test-cuda-projection-fixture", "Projection fixture dispatch controls with CPU-only CUDA calls"); 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"), @@ -251,6 +252,15 @@ pub fn hostTests(b: *std.Build, draft_ids: *std.Build.Module, build_options: *st }) })); step.dependOn(&reference.step); reference_step.dependOn(&reference.step); + const projection = b.addRunArtifact(b.addTest(.{ .root_module = b.createModule(.{ + .root_source_file = b.path("zig/tests/cuda/projection_f32_host_tests.zig"), + .target = host, + .optimize = mode, + .link_libc = true, + .imports = &.{ .{ .name = "cuda", .module = cuda }, .{ .name = "core", .module = mods.core } }, + }) })); + step.dependOn(&projection.step); + projection_step.dependOn(&projection.step); } const native_modules = engines(b, host, .debug, cuda, mods.core, mods.lanes, mods.nemotron, mods.heat); const native = native_modules.engines; diff --git a/zig/tests/cuda/projection_f32_host_tests.zig b/zig/tests/cuda/projection_f32_host_tests.zig new file mode 100644 index 0000000000..5d38e2e29c --- /dev/null +++ b/zig/tests/cuda/projection_f32_host_tests.zig @@ -0,0 +1,109 @@ +//! Exercise the fixture's actual poison-and-dispatch path with CPU buffers and no-op CUDA launches. +const std = @import("std"); +const cuda = @import("cuda"); +const fixture = @import("projection_f32_tests.zig"); +const Rig = @import("quant_ref_tests.zig").RigT; + +const Mock = struct { + var destination: []u8 = &.{}; + var scratch: []u8 = &.{}; + var fills: usize = 0; + var launches: usize = 0; + var fresh = true; + + fn fill(ptr: cuda.abi.DevicePtr, byte: u8, n: usize, _: cuda.abi.Stream) callconv(.c) cuda.abi.Result { + const bytes: [*]u8 = @ptrFromInt(ptr); + @memset(bytes[0..n], byte); + fills += 1; + return cuda.abi.success; + } + + fn observe() cuda.abi.Result { + for (destination) |byte| fresh = fresh and byte == 0xff; + for (scratch) |byte| fresh = fresh and byte == 0xff; + launches += 1; + return cuda.abi.success; + } + + fn plain(_: cuda.abi.Function, _: c_uint, _: c_uint, _: c_uint, _: c_uint, _: c_uint, _: c_uint, _: c_uint, _: cuda.abi.Stream, _: ?[*]?*anyopaque, _: ?[*]?*anyopaque) callconv(.c) cuda.abi.Result { + return observe(); + } + + fn extended(_: *const cuda.abi.LaunchConfig, _: cuda.abi.Function, _: ?[*]?*anyopaque, _: ?[*]?*anyopaque) callconv(.c) cuda.abi.Result { + return observe(); + } + + fn sync(_: cuda.abi.Stream) callconv(.c) cuda.abi.Result { + 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]); + return cuda.abi.success; + } +}; + +test "every F32 and BF16 fixture dispatch rejects a no-op after a valid stale prefix" { + const a = std.testing.allocator; + var driver: cuda.Driver = undefined; + driver.api.cuMemsetD8Async = Mock.fill; + driver.api.cuLaunchKernel = Mock.plain; + driver.api.cuLaunchKernelEx = Mock.extended; + driver.api.cuStreamSynchronize = Mock.sync; + driver.api.cuMemcpyDtoH_v2 = Mock.copy; + const stream: cuda.Stream = .{ .d = &driver, .handle = @ptrFromInt(1) }; + const f: cuda.Function = .{ .d = &driver, .handle = null }; + var lane: cuda.qmmf.Lane = .{ + .mod = .{ .d = &driver, .handle = null }, + .fns = @splat(@splat(f)), + .f32_fns = @splat(@splat(f)), + .reduce = @splat(f), + .clusters = false, + }; + const weight: cuda.qmmf.Weight = .{ .mode = .fp4, .codes = 0, .scales = 0, .n = 70, .k = 1024, .npad = 128 }; + const elements = 256 * weight.n; + const expected = try a.alloc(f64, elements); + defer a.free(expected); + const tolerance = try a.alloc(f64, elements); + defer a.free(tolerance); + @memset(expected, 0); + @memset(tolerance, 1); + var arena = std.heap.ArenaAllocator.init(a); + defer arena.deinit(); + var rig: Rig = .{ + .gpu = .{ .d = &driver, .ctx = undefined, .gpa = a, .io = std.testing.io }, + .a = arena.allocator(), + .s = stream, + .rng = undefined, + }; + for ([_]cuda.qmmf.Out{ .f32, .bf16 }) |format| { + const width: usize = if (format == .f32) 4 else 2; + const destination = try a.alloc(u8, elements * width); + defer a.free(destination); + const scratch = try a.alloc(u8, 8 * elements * 4); + defer a.free(scratch); + const out: cuda.DeviceBuffer = .{ .d = &driver, .ptr = @intFromPtr(destination.ptr), .len = destination.len }; + const part: cuda.DeviceBuffer = .{ .d = &driver, .ptr = @intFromPtr(scratch.ptr), .len = scratch.len }; + Mock.destination = destination; + Mock.scratch = scratch; + for ([_]bool{ true, false }) |clusters| { + lane.clusters = clusters; + for ([_]usize{ 256, 1, 7, 17, 33 }) |rows| { + @memset(destination, 0); + @memset(scratch, 0); + const count = rows * weight.n; + try rig.close("valid stale prefix", out.ptr, format == .f32, expected[0..count], tolerance[0..count]); + Mock.fills = 0; + Mock.launches = 0; + Mock.fresh = true; + try fixture.dispatch(&lane, stream, 0, 1088, rows, weight, out, part, format); + try std.testing.expectEqual(@as(usize, 2), Mock.fills); + try std.testing.expect(Mock.launches > 0); + try std.testing.expect(Mock.fresh); + try std.testing.expectError(error.TestFailed, rig.close("no-op projection", out.ptr, format == .f32, expected[0..count], tolerance[0..count])); + } + } + } +} diff --git a/zig/tests/cuda/projection_f32_tests.zig b/zig/tests/cuda/projection_f32_tests.zig index 19a62fdf4f..3edd66a243 100644 --- a/zig/tests/cuda/projection_f32_tests.zig +++ b/zig/tests/cuda/projection_f32_tests.zig @@ -9,6 +9,13 @@ const bf = reference.helpers.bfOf; const bf16 = reference.helpers.bf16Of; const Mode = cuda.qmmf.Mode; +/// One fixture dispatch starts with fresh destination and reduction poison, including a reused row prefix. +pub fn dispatch(lane: *const cuda.qmmf.Lane, s: cuda.Stream, x: u64, ldx: usize, rows: usize, w: cuda.qmmf.Weight, out: cuda.DeviceBuffer, part: cuda.DeviceBuffer, format: cuda.qmmf.Out) !void { + try out.fill8(0xff, s.handle); + try part.fill8(0xff, s.handle); + try lane.matmulOut(s, x, ldx, rows, w, out.ptr, format, part.ptr); +} + pub fn run(gpu: check.Gpu) !void { var arena = std.heap.ArenaAllocator.init(gpu.gpa); defer arena.deinit(); @@ -88,7 +95,6 @@ fn projection(r: *Rig, lane: *cuda.qmmf.Lane, mode: Mode) !void { defer part_buf.free(); const out = out_buf.ptr; const out_bf16 = bf16_buf.ptr; - const part = part_buf.ptr; const want = try a.alloc(f64, m * n); const tol = try a.alloc(f64, m * n); for (0..m) |row| for (0..n) |col| { @@ -117,15 +123,11 @@ fn projection(r: *Rig, lane: *cuda.qmmf.Lane, mode: Mode) !void { for ([_]bool{ true, false }) |clusters| { lane.clusters = clusters and (try r.gpu.ctx.capability()) >= 90; for ([_]usize{ 256, 1, 7, 17, 33 }) |rows| { - try out_buf.fill8(0xff, r.s.handle); - try part_buf.fill8(0xff, r.s.handle); - try lane.matmulF32(r.s, dx, ldx, rows, w, out, part); + try dispatch(lane, r.s, dx, ldx, rows, w, out_buf, part_buf, .f32); try r.close("FP32 projection", out, true, want[0 .. rows * n], tol[0 .. rows * n]); const got = try r.back(out, rows * n * 4); if (baseline) |b| try check.sameBytes("FP32 row/reduction invariance", got, b[0..got.len]) else baseline = got; - try bf16_buf.fill8(0xff, r.s.handle); - try part_buf.fill8(0xff, r.s.handle); - try lane.matmul(r.s, dx, ldx, rows, w, out_bf16, part); + try dispatch(lane, r.s, dx, ldx, rows, w, bf16_buf, part_buf, .bf16); const rounded = try r.back(out_bf16, rows * n * 2); for (0..rows * n) |i| { const v = std.mem.readInt(u32, got[i * 4 ..][0..4], .little);