Fix CUDA backend merge resolution truncations

Five additional CUDA-specific issues from the union-style PrismML merge
resolution surfaced during nvcc compilation (CPU build skipped these
files entirely):

- common.cuh: removed duplicate ggml_cuda_type_traits<GGML_TYPE_Q1_0>
  specialization (line 928 was redundant with line 977)
- dequantize.cuh: removed duplicated body inside dequantize_q1_0_g128
  (caused "already declared in current scope" for bit_index_0 etc.)
- mmq.cu: removed duplicate `case GGML_TYPE_Q1_0:` in
  ggml_cuda_should_use_mmq switch
- mmq.cuh: added missing closing brace after load_tiles_q1_0_g128
  before load_tiles_q2_0 began
- vecdotq.cuh: re-added vec_dot_q2_0_q8_1 function (dropped by union
  resolver, but referenced in mmvq.cu switch)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

Author: retroheim

ggml/src/ggml-cuda/common.cuh

@@ -973,13 +973,6 @@ struct ggml_cuda_type_traits<GGML_TYPE_Q8_0> {
     static constexpr int qi = QI8_0;
 };
 
-template<>
-struct ggml_cuda_type_traits<GGML_TYPE_Q1_0> {
-    static constexpr int qk = QK1_0;
-    static constexpr int qr = QR1_0;
-    static constexpr int qi = QI1_0;
-};
-
 template<>
 struct ggml_cuda_type_traits<GGML_TYPE_Q1_0_g128> {
     static constexpr int qk = QK1_0_g128;

ggml/src/ggml-cuda/dequantize.cuh

@@ -44,22 +44,6 @@ static __device__ __forceinline__ void dequantize_q1_0_g128(const void * vx, con
 
     v.x = bit_0 ? d : neg_d;
     v.y = bit_1 ? d : neg_d;
-
-    const int bit_index_0 = iqs;
-    const int bit_index_1 = iqs + 1;
-
-    const int byte_index_0 = bit_index_0 / 8;
-    const int bit_offset_0 = bit_index_0 % 8;
-
-    const int byte_index_1 = bit_index_1 / 8;
-    const int bit_offset_1 = bit_index_1 % 8;
-
-    // Extract bits: 1 = +d, 0 = -d (branchless)
-    const int bit_0 = (x[ib].qs[byte_index_0] >> bit_offset_0) & 1;
-    const int bit_1 = (x[ib].qs[byte_index_1] >> bit_offset_1) & 1;
-
-    v.x = (2*bit_0 - 1) * d;
-    v.y = (2*bit_1 - 1) * d;
 }
 
 static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const int64_t ib, const int iqs, float2 & v){

ggml/src/ggml-cuda/mmq.cu

@@ -283,7 +283,6 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
         // TODO: Q1_0 and Q1_0_g128 MMQ implementation exists but is currently disabled due to accuracy issues
         case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q1_0_g128:
-        case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q2_0:
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:

ggml/src/ggml-cuda/mmq.cuh

@@ -455,6 +455,8 @@ template <int mmq_y, bool need_check> static __device__ __forceinline__ void loa
         x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + ksx] = bxi->d;
     }
 #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)
+}
+
 template <int mmq_y, bool need_check> static __device__ __forceinline__ void load_tiles_q2_0(
     const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
     constexpr int nwarps = mmq_get_nwarps_device();

ggml/src/ggml-cuda/vecdotq.cuh

@@ -788,6 +788,54 @@ static __device__ __forceinline__ float vec_dot_q1_0_g128_q8_1(
     return d1 * ds8f.x * sumi;
 }
 
+static __device__ __forceinline__ float vec_dot_q2_0_q8_1(
+    const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
+
+    const block_q2_0 * bq2_0 = (const block_q2_0 *) vbq + kbx;
+
+    const float d2 = bq2_0->d;
+
+    const block_q8_1 * bq8_1_chunk = bq8_1 + iqs;
+
+    const int offset = iqs * 8;
+    const int v0 = bq2_0->qs[offset + 0] | (bq2_0->qs[offset + 1] << 8) |
+                   (bq2_0->qs[offset + 2] << 16) | (bq2_0->qs[offset + 3] << 24);
+    const int v1 = bq2_0->qs[offset + 4] | (bq2_0->qs[offset + 5] << 8) |
+                   (bq2_0->qs[offset + 6] << 16) | (bq2_0->qs[offset + 7] << 24);
+
+    int vi_bytes[8];
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        const int shift = j * 8;
+        const int codes = (v0 >> shift) & 0xFF;
+        const int c0 = ((codes >> 0) & 0x3) - 1;
+        const int c1 = ((codes >> 2) & 0x3) - 1;
+        const int c2 = ((codes >> 4) & 0x3) - 1;
+        const int c3 = ((codes >> 6) & 0x3) - 1;
+        vi_bytes[j] = (c0 & 0xFF) | ((c1 & 0xFF) << 8) | ((c2 & 0xFF) << 16) | ((c3 & 0xFF) << 24);
+    }
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        const int shift = j * 8;
+        const int codes = (v1 >> shift) & 0xFF;
+        const int c0 = ((codes >> 0) & 0x3) - 1;
+        const int c1 = ((codes >> 2) & 0x3) - 1;
+        const int c2 = ((codes >> 4) & 0x3) - 1;
+        const int c3 = ((codes >> 6) & 0x3) - 1;
+        vi_bytes[4 + j] = (c0 & 0xFF) | ((c1 & 0xFF) << 8) | ((c2 & 0xFF) << 16) | ((c3 & 0xFF) << 24);
+    }
+
+    int sumi = 0;
+#pragma unroll
+    for (int j = 0; j < 8; ++j) {
+        const int u = get_int_b4(bq8_1_chunk->qs, j);
+        sumi = ggml_cuda_dp4a(vi_bytes[j], u, sumi);
+    }
+
+    const float d8 = __low2float(bq8_1_chunk->ds);
+    return d2 * d8 * sumi;
+}
+
 static __device__ __forceinline__ float vec_dot_q4_0_q8_1(
     const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {