From d4e21f0d34d180e22b1a11d3d3ac1675f3f52cd5 Mon Sep 17 00:00:00 2001 From: Chris Chuter Date: Thu, 14 May 2026 12:14:11 -0500 Subject: [PATCH] cuda: software-FP16 fallback for dequantize_V_q{4,5,8}_* templates MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Clean builds of feat/v4-port-cuda fail with 5 static_assert errors in fattn-common.cuh when the CMake CUDA arch matrix includes pre-Pascal archs (50-virtual is in the default matrix). The five dequantize_V_q* templates have an `#ifdef FP16_AVAILABLE if constexpr (T==half) ... else` structure; when FP16_AVAILABLE is undefined the preprocessor strips the half arm and `T=half` falls into the trailing `static_assert(..., "bad type")` else. The instantiation `get_dequantize_V()` is reached unconditionally from lightning_indexer_kernel_vec for quantized K-cache types, and that kernel is dispatched on all CCs (the wmma path is the only Ampere+ special-case; everything else falls through). So a runtime NO_DEVICE_CODE trap is not acceptable — we need a real software-FP16 path. This commit adds an `#else` arm in each of the five templates that: 1. converts the per-block scale (and min) to float via __half2float / __half22float2 (both available on all supported CCs); 2. computes the dequantized scalars in float; 3. packs adjacent pairs into half2 via __floats2half2_rn (RTNE conversion, available on all supported CCs). The FP16-available branch is byte-identical to before. The new branch produces results at strictly higher intermediate precision than the FP16-hardware path (FP32 fmadd then RTNE-to-half, vs FP16 mul/add), matching the established pattern used elsewhere in common.cuh (e.g. ggml_cuda_mad fallbacks). Latent regression introduced by upstream-merge 860b585ea; previously hidden by stale incremental builds. Co-Authored-By: Claude Opus 4.7 (1M context) --- ggml/src/ggml-cuda/fattn-common.cuh | 52 +++++++++++++++++++++++++++++ 1 file changed, 52 insertions(+) diff --git a/ggml/src/ggml-cuda/fattn-common.cuh b/ggml/src/ggml-cuda/fattn-common.cuh index 1cd65ad88d41..0991197f192e 100644 --- a/ggml/src/ggml-cuda/fattn-common.cuh +++ b/ggml/src/ggml-cuda/fattn-common.cuh @@ -400,6 +400,18 @@ static __device__ __forceinline__ void dequantize_V_q4_0(const void * __restrict ((half2 *) dst)[l0/2] = d * make_half2(q8[l0 + 0], q8[l0 + 1]); } } else +#else // FP16_AVAILABLE + // Software-FP16 fallback: compute in float, pack via __floats2half2_rn (RTNE). + // Required because lightning_indexer_kernel_vec instantiates for the + // full CC matrix (50/61/70/...), but FP16_AVAILABLE is undefined for CC < 600. + if constexpr (std::is_same_v) { + const float d = __half2float(x[ib].d); + +#pragma unroll + for (int l0 = 0; l0 < ne; l0 += 2) { + ((half2 *) dst)[l0/2] = __floats2half2_rn(d * q8[l0 + 0], d * q8[l0 + 1]); + } + } else #endif // FP16_AVAILABLE if constexpr (std::is_same_v) { const float d = x[ib].d; @@ -440,6 +452,16 @@ static __device__ __forceinline__ void dequantize_V_q4_1(const void * __restrict ((half2 *) dst)[l0/2] = d * make_half2(q8[l0 + 0], q8[l0 + 1]) + m; } } else +#else // FP16_AVAILABLE + // Software-FP16 fallback: see dequantize_V_q4_0 for the rationale. + if constexpr (std::is_same_v) { + const float2 dm = __half22float2(x[ib].dm); + +#pragma unroll + for (int l0 = 0; l0 < ne; l0 += 2) { + ((half2 *) dst)[l0/2] = __floats2half2_rn(dm.x * q8[l0 + 0] + dm.y, dm.x * q8[l0 + 1] + dm.y); + } + } else #endif // FP16_AVAILABLE if constexpr (std::is_same_v) { const float2 dm = __half22float2(x[ib].dm); @@ -490,6 +512,16 @@ static __device__ __forceinline__ void dequantize_V_q5_0(const void * __restrict ((half2 *) dst)[l0/2] = d * make_half2(q8[l0 + 0], q8[l0 + 1]); } } else +#else // FP16_AVAILABLE + // Software-FP16 fallback: see dequantize_V_q4_0 for the rationale. + if constexpr (std::is_same_v) { + const float d = __half2float(x[ib].d); + +#pragma unroll + for (int l0 = 0; l0 < ne; l0 += 2) { + ((half2 *) dst)[l0/2] = __floats2half2_rn(d * q8[l0 + 0], d * q8[l0 + 1]); + } + } else #endif // FP16_AVAILABLE if constexpr (std::is_same_v) { const float d = x[ib].d; @@ -540,6 +572,16 @@ static __device__ __forceinline__ void dequantize_V_q5_1(const void * __restrict ((half2 *) dst)[l0/2] = d * make_half2(q8[l0 + 0], q8[l0 + 1]) + m; } } else +#else // FP16_AVAILABLE + // Software-FP16 fallback: see dequantize_V_q4_0 for the rationale. + if constexpr (std::is_same_v) { + const float2 dm = __half22float2(x[ib].dm); + +#pragma unroll + for (int l0 = 0; l0 < ne; l0 += 2) { + ((half2 *) dst)[l0/2] = __floats2half2_rn(dm.x * q8[l0 + 0] + dm.y, dm.x * q8[l0 + 1] + dm.y); + } + } else #endif // FP16_AVAILABLE if constexpr (std::is_same_v) { const float2 dm = __half22float2(x[ib].dm); @@ -573,6 +615,16 @@ static __device__ __forceinline__ void dequantize_V_q8_0(const void * __restrict ((half2 *) dst)[l0/2] = d * make_half2(qs[l0 + 0], qs[l0 + 1]); } } else +#else // FP16_AVAILABLE + // Software-FP16 fallback: see dequantize_V_q4_0 for the rationale. + if constexpr (std::is_same::value) { + const float d = __half2float(x[ib].d); + +#pragma unroll + for (int l0 = 0; l0 < ne; l0 += 2) { + ((half2 *) dst)[l0/2] = __floats2half2_rn(d * qs[l0 + 0], d * qs[l0 + 1]); + } + } else #endif // FP16_AVAILABLE if constexpr (std::is_same::value) { const float d = x[ib].d;