From 872057bd79b0ec52ef0b078823a744112c0a001e Mon Sep 17 00:00:00 2001 From: Victor Cruz Date: Mon, 24 Aug 2026 02:24:37 -0700 Subject: [PATCH] fix(gguf): unsupported quant types returned uninitialized memory None of the five switch (type) blocks in gguf_kernel.cu had a default: case, and the output tensor is allocated with torch::empty, so an unsupported quant type returned uninitialized memory rather than raising. ggml_moe_get_block_size fell through to return 0, which then divides or sizes a launch by zero downstream. The gap is reachable: ggml_mul_mat_vec_a8 and ggml_moe_a8_vec handle 19 types while ggml_mul_mat_a8 and ggml_moe_a8 handle 10, because the I-quants have no MMQ kernel. Any caller routing an I-quant to the MMQ path gets silent garbage, and the only symptom is a model that loads, runs at full speed, and produces fluent nonsense. Each default: now raises via TORCH_CHECK naming the function and the type families that entry point supports. ggml_dequantize also null-checks the ggml_get_to_cuda function pointer, which returns nullptr for an unknown type and was being called unconditionally. No kernel math changed. A comment above the first guard notes these kernels are vendored from sgl-kernel so a future re-vendor does not drop them. --- .../freetoken/kernel/csrc/gguf/gguf_kernel.cu | 28 ++++++++++++++++++- 1 file changed, 27 insertions(+), 1 deletion(-) diff --git a/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu b/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu index d88960d5..db210646 100644 --- a/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu +++ b/python/freetoken/kernel/csrc/gguf/gguf_kernel.cu @@ -83,8 +83,14 @@ torch::Tensor ggml_dequantize( at::Tensor DW = torch::empty({m, n}, options); cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream(); + // These kernels are vendored from sgl-kernel/llama.cpp; the guards below are a FreeToken addition + // to prevent silent data corruption from unsupported quant types. DISPATCH_FLOAT_TYPES(DW.scalar_type(), "ggml_dequantize", [&] { auto to_cuda = ggml_get_to_cuda(type); + TORCH_CHECK(to_cuda != nullptr, + "ggml_dequantize: unsupported GGUF quant type ", type, + " (dequant kernels exist for Q4_0/Q4_1/Q5_0/Q5_1/Q8_0/Q2_K-Q6_K/IQ2_XXS/" + "IQ2_XS/IQ3_XXS/IQ1_S/IQ4_NL/IQ3_S/IQ2_S/IQ4_XS/IQ1_M)"); to_cuda((void*)W.data_ptr(), (scalar_t*)DW.data_ptr(), m * n, stream); }); @@ -184,6 +190,10 @@ torch::Tensor ggml_mul_mat_vec_a8( mul_mat_vec_iq1_m_q8_1_cuda( (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (scalar_t*)Y.data_ptr(), col, row, vecs, stream); break; + default: + TORCH_CHECK(false, "ggml_mul_mat_vec_a8: unsupported GGUF quant type ", type, + " (MMVQ kernels exist for Q4_0/Q4_1/Q5_0/Q5_1/Q8_0/Q2_K-Q6_K/IQ2_XXS/IQ2_XS/" + "IQ3_XXS/IQ1_S/IQ4_NL/IQ3_S/IQ2_S/IQ4_XS/IQ1_M)"); } }); return Y; @@ -327,6 +337,10 @@ torch::Tensor ggml_mul_mat_a8( row, stream); break; + default: + TORCH_CHECK(false, "ggml_mul_mat_a8: unsupported GGUF quant type ", type, + " (MMQ kernels exist only for Q4_0/Q4_1/Q5_0/Q5_1/Q8_0/Q2_K-Q6_K; " + "I-quants must route through ggml_dequantize)"); } }); return Y; @@ -533,6 +547,10 @@ torch::Tensor ggml_moe_a8( sorted_token_ids.sizes()[0], stream); break; + default: + TORCH_CHECK(false, "ggml_moe_a8: unsupported GGUF quant type ", type, + " (MMQ kernels exist only for Q4_0/Q4_1/Q5_0/Q5_1/Q8_0/Q2_K-Q6_K; " + "I-quants must route through ggml_dequantize)"); } }); return Y; @@ -804,6 +822,10 @@ torch::Tensor ggml_moe_a8_vec( quant_X.stride(0), stream); break; + default: + TORCH_CHECK(false, "ggml_moe_a8_vec: unsupported GGUF quant type ", type, + " (MMVQ kernels exist for Q4_0/Q4_1/Q5_0/Q5_1/Q8_0/Q2_K-Q6_K/IQ2_XXS/IQ2_XS/" + "IQ3_XXS/IQ1_S/IQ4_NL/IQ3_S/IQ2_S/IQ4_XS/IQ1_M)"); } }); return Y; @@ -831,8 +853,12 @@ int64_t ggml_moe_get_block_size(int64_t type) { return MOE_X_Q5_K; case 14: return MOE_X_Q6_K; + default: + TORCH_CHECK(false, "ggml_moe_get_block_size: unsupported GGUF quant type ", type, + " (MMQ kernels exist only for Q4_0/Q4_1/Q5_0/Q5_1/Q8_0/Q2_K-Q6_K; " + "I-quants must route through ggml_dequantize)"); + return 0; // unreachable but silences compiler warning } - return 0; } // ---- FreeToken pybind bindings (donor registers these via TORCH_LIBRARY; we