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/build/cuda.zig b/zig/build/cuda.zig index 684d6c3cc9..ee6685bf9d 100644 --- a/zig/build/cuda.zig +++ b/zig/build/cuda.zig @@ -240,6 +240,28 @@ 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)"); + 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"), + .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 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; // the HTTP server's unit tests (zig/src/server/root.zig), which the macOS build runs, on Linux hosts too diff --git a/zig/kernels/cuda/qmmf.cu b/zig/kernels/cuda/qmmf.cu index 65cf08fd22..6ef58c2406 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); @@ -343,6 +356,42 @@ 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 reduce_kernel(const float* __restrict__, void* __restrict__, long long, int, float); +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/core/contract/registry.zig b/zig/src/core/contract/registry.zig index b79337b5c5..d1e2bf0307 100644 --- a/zig/src/core/contract/registry.zig +++ b/zig/src/core/contract/registry.zig @@ -20,23 +20,59 @@ 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 } } }, + .{ .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 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 } } }, 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 dbb86bb8d0..a5031b6bcd 100644 --- a/zig/src/cuda/qlinear.zig +++ b/zig/src/cuda/qlinear.zig @@ -34,6 +34,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 @@ -41,14 +42,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, }; } }; @@ -81,6 +82,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 if (q.act != 0 and l.ck != null) l.checkpoint(s, l.ck.?, q, in, out, rows) else l.onLane(s, q, .fp8, in, out, rows), .bf16 => |q| return l.onLane(s, q, .bf16, in, out, rows), @@ -96,6 +98,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 if (q.act != 0 and l.ck != null) l.checkpoint(s, l.ck.?, q, in, out, rows) else l.onPrompt(s, q, .fp8, in, out, rows), .bf16 => |q| return l.onLane(s, q, .bf16, in, out, rows), @@ -385,6 +388,8 @@ 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)); const f8: qmmf.Weight = .{ .mode = .fp8g, .codes = 0, .scales = 0, .n = 64, .k = 64, .npad = 128 }; try std.testing.expectError(error.EpilogueNotBuilt, l.decodeF32(s, .{ .fp8g = f8 }, .{ .x = 0 }, 0, 1)); 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 ec7fcd6e04..8fb23389ad 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); @@ -85,7 +85,7 @@ pub const reduce_symbols = [2][:0]const u8{ }; const row_tiles = [_]u32{ 16, 32, 64 }; -/// What a call writes: bf16 rows, or the fp32 sums unrounded (FP4 only, a tensor-parallel rank's partial rows). +/// What a call writes: bf16 rows, or the fp32 sums unrounded (a tensor-parallel rank's partial rows). pub const Out = enum { bf16, f32 }; fn loadTiles(mod: Module, comptime mode: Mode, comptime wide: bool) ![7]Function { @@ -103,7 +103,7 @@ fn loadTiles(mod: Module, comptime mode: Mode, comptime wide: bool) ![7]Function pub const Lane = struct { mod: Module, fns: [modes.len][7]Function, // per mode: lane16, lane16c, lane32, lane32c, lane64, lane64c, fused - fp4_f32: [7]Function, // FP4's tiles writing fp32 + f32_fns: [modes.len][7]Function, reduce: [2]Function, // by Out clusters: bool, // sm_90 on: slices meet in a cluster's shared memory @@ -114,11 +114,14 @@ pub const Lane = struct { var l: Lane = .{ .mod = mod, .fns = undefined, - .fp4_f32 = try loadTiles(mod, .fp4, true), + .f32_fns = undefined, .reduce = undefined, .clusters = capability_major >= 9, }; - inline for (modes, 0..) |mode, mi| l.fns[mi] = try loadTiles(mod, mode, false); + inline for (modes, 0..) |mode, mi| { + l.fns[mi] = try loadTiles(mod, mode, false); + l.f32_fns[mi] = try loadTiles(mod, mode, true); + } for (&l.reduce, reduce_symbols) |*f, name| f.* = try mod.function(name); return l; } @@ -138,11 +141,15 @@ pub const Lane = struct { return l.matmulOut(s, x, ldx, m, w, y, .bf16, part); } - /// matmul() writing y as `out` says: bf16, or fp32 (FP4 only). + /// Preserve 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.matmulOut(s, x, ldx, m, w, y, .f32, part); + } + + /// matmul() writing y as `out` says: bf16, or fp32. pub fn matmulOut(l: Lane, s: Stream, x: u64, ldx: usize, m: usize, w: Weight, y: u64, out: Out, part: u64) !void { - if (out == .f32 and w.mode != .fp4) return error.EpilogueNotBuilt; const p = plan(m, w.n, w.k, l.clusters); - const set = if (out == .f32) l.fp4_f32 else l.fns[modeIndex(w.mode)]; + const set = (if (out == .f32) 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; @@ -220,6 +227,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)); @@ -243,16 +253,9 @@ 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_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), - ); - try std.testing.expectEqualStrings( - "_ZN7tf_qmmf11qmmf_kernelILi0ELi32ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", - symbol(.fp4, 32, true, 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)); + try std.testing.expectEqualStrings("_ZN7tf_qmmf11qmmf_kernelILi0ELi32ELi64ELi1ELi4ELi4ELb1ELb1ELb0EEEvPK13__nv_bfloat16PKhS5_fPvPfiiiiiii", symbol(.fp4, 32, true, false, true)); } diff --git a/zig/src/cuda/root.zig b/zig/src/cuda/root.zig index 6c4e0991b6..f49c4de197 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"); @@ -58,6 +59,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/contract_mxfp8.zig b/zig/tests/cuda/contract_mxfp8.zig new file mode 100644 index 0000000000..278cb78c4c --- /dev/null +++ b/zig/tests/cuda/contract_mxfp8.zig @@ -0,0 +1,115 @@ +//! MXFP8 projections through the common non-finite, repeat, shape and accumulator gates. +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; + +fn depth(shape: c.Shape) usize { + return shape.k / 16 + 24; +} + +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 }, + .{ .rows = 7, .n = 13952, .k = 2560 }, +}; + +pub const Product = struct { + rig: Rig, + lane: *const cuda.qmmf.Lane, + 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 { + return .{ .spec = .{ + .name = if (self.clusters) "cuda MXFP8 lane clustered slices" else "cuda MXFP8 lane explicit slice reduction", + .accum = .f32, + .depth = depth, + .out = .bf16, + .tiles = .{ .{ .size = 16 }, .{ .size = 64 }, .{ .size = 512, .step = 64, .min = 64 } }, + .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 % 64 != 0) return false; + const a = self.rig.gpu.gpa; + const d = try Dense.init(a, shape, 1); + self.d = d; + d.fillX(fill, seed, 16); + var prng = std.Random.DefaultPrng.init(seed ^ 0x8f); + const random = prng.random(); + const n = shape.n; + const k = shape.k; + const codes = try a.alloc(u8, n * k); + defer a.free(codes); + const scales = try a.alloc(u8, n * k / 32); + defer a.free(scales); + for (scales) |*s| s.* = if (fill == .stress) 127 else 123 + random.uintLessThan(u8, 8); + for (0..n) |j| for (0..k) |i| { + codes[j * k + i] = if (fill == .stress) 0x38 + random.uintLessThan(u8, 24) else random.uintLessThan(u8, 0x7f); + if (fill != .stress and random.boolean()) codes[j * k + i] |= 0x80; + }; + if (fill == .stress) for (0..n) |j| for (k / 2..k) |i| { + codes[j * k + i] = codes[j * k + i - k / 2] | 0x80; + }; + for (0..n) |j| for (0..k) |i| { + const scale = std.math.ldexp(@as(f64, 1), @as(i32, scales[j * (k / 32) + i / 32]) - 127); + d.w[j * k + i] = core.quant.fp8.e4m3(codes[j * k + i]) * scale; + }; + const npad = cuda.mxfp8.padded(n); + const tiled_codes = try a.alloc(u8, npad * k); + defer a.free(tiled_codes); + const packed_scales = try a.alloc(u8, npad * k / 32); + defer a.free(packed_scales); + cuda.mxfp8.packCodes(tiled_codes, codes, n, k); + cuda.mxfp8.packScales(packed_scales, scales, n, k); + self.weight = cuda.mxfp8.weight(try self.rig.input(tiled_codes), try self.rig.input(packed_scales), n, k); + 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 * n * 2); + return true; + } + + pub fn run(self: *Product) !void { + const d = self.d.?; + try self.lane.matmul(self.rig.s, self.x, d.shape.k, d.shape.rows, self.weight, self.out, 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_projection_f32.zig b/zig/tests/cuda/contract_projection_f32.zig new file mode 100644 index 0000000000..1e172dec30 --- /dev/null +++ b/zig/tests/cuda/contract_projection_f32.zig @@ -0,0 +1,159 @@ +//! 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 }, + .{ .rows = 7, .n = 2560, .k = 10240 }, + .{ .rows = 7, .n = 13952, .k = 2560 }, +}; + +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 e87a2e90e5..e1f58336a2 100644 --- a/zig/tests/cuda/contract_tests.zig +++ b/zig/tests/cuda/contract_tests.zig @@ -7,6 +7,8 @@ const check = @import("check.zig"); 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; @@ -86,6 +88,19 @@ pub fn run(gpu: check.Gpu) !void { } defer for (&mm) |*m| m.deinit(); for (&mm) |*m| try kernels.append(a, m.kernel()); + var mx = [_]mx_k.Product{ + .{ .rig = .{ .gpu = gpu, .s = s }, .lane = &lane, .clusters = true }, + .{ .rig = .{ .gpu = gpu, .s = s }, .lane = &lane_reduce, .clusters = false }, + }; + 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()); diff --git a/zig/tests/cuda/main.zig b/zig/tests/cuda/main.zig index fe98682c86..3a1944118f 100644 --- a/zig/tests/cuda/main.zig +++ b/zig/tests/cuda/main.zig @@ -17,6 +17,8 @@ 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 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"); @@ -80,6 +82,8 @@ 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) \\ nvfp4 NVFP4 projections against recorded lane-matmul and prompt-GEMM fixture bytes @@ -155,6 +159,8 @@ 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, "glm-ref")) return glm_ref_tests.run(gpu); if (std.mem.eql(u8, cmd, "modelopt")) return modelopt_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", .{}); +} 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 new file mode 100644 index 0000000000..3edd66a243 --- /dev/null +++ b/zig/tests/cuda/projection_f32_tests.zig @@ -0,0 +1,139 @@ +//! 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; + +/// 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(); + 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)); + 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 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 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 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); + 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}); +} 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;