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 <half, 4> for the
+    // full CC matrix (50/61/70/...), but FP16_AVAILABLE is undefined for CC < 600.
+    if constexpr (std::is_same_v<T, half>) {
+        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<T, float>) {
         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<T, half>) {
+        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<T, float>) {
         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<T, half>) {
+        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<T, float>) {
         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<T, half>) {
+        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<T, float>) {
         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<T, half>::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<T, float>::value) {
         const float d = x[ib].d;