tests : add DT3 GPU vs CPU parity test

Checks the GPU backend against the validated CPU path: GET_ROWS
dequantization must match dequantize_row_dt3 bit by bit on directed
blocks (region boundaries 79/80 and 119/120, qh elements, negative
scales) and on raw random bytes; MUL_MAT must match a double precision
reference from the dequantized weights, with activations whose q8_1
quantization is exact, plus a manual trit-sum check with non-trivial qh.
Skips cleanly when no GPU backend is available.

ggml/src/ggml-common.h

@@ -99,6 +99,9 @@ typedef sycl::half2 ggml_half2;
 #define QI2_0 (QK2_0 / 32)
 #define QR2_0 1
 
+#define QI_DT3 (QK_DT3 / 32)
+#define QR_DT3 1
+
 
 #define QI4_0 (QK4_0 / (4 * QR4_0))
 #define QR4_0 2

ggml/src/ggml-cuda/common.cuh

@@ -962,6 +962,26 @@ static __device__ __forceinline__ float get_alibi_slope(
     return powf(base, exph);
 }
 
+// decode element i (0..QK_DT3) of one packed DT3 ternary plane to a trit in {-1, 0, +1}
+// layout per plane (see block_dt3): qs[0..16) hold elements m + n*16 (m = 0..16, n = 0..5),
+// qs[16..24) hold elements 80 + m + n*8 (m = 0..8, n = 0..5), qh[0..2) hold elements
+// 120 + j + n*2 (j = 0..2, n = 0..4) — a qh byte stores only 4 trits, its 5th base-3
+// digit is packing padding that always decodes to -1 and must never be read
+static __device__ __forceinline__ int ggml_cuda_dt3_get_trit(
+        const uint8_t * __restrict__ qs, const uint8_t * __restrict__ qh, const int i) {
+    const uint8_t pow3[5] = {1, 3, 9, 27, 81};
+
+    uint8_t q; // the multiplications below wrap around in uint8_t on purpose
+    if (i < 80) {
+        q = qs[i % 16] * pow3[i / 16];
+    } else if (i < 120) {
+        q = qs[16 + (i - 80) % 8] * pow3[(i - 80) / 8];
+    } else {
+        q = qh[(i - 120) % 2] * pow3[(i - 120) / 2];
+    }
+    return (int) (((uint16_t) q * 3) >> 8) - 1;
+}
+
 template <ggml_type type>
 struct ggml_cuda_type_traits;
 
@@ -985,6 +1005,13 @@ struct ggml_cuda_type_traits<GGML_TYPE_Q2_0> {
     static constexpr int qi = QI2_0;
 };
 
+template<>
+struct ggml_cuda_type_traits<GGML_TYPE_DT3> {
+    static constexpr int qk = QK_DT3;
+    static constexpr int qr = QR_DT3;
+    static constexpr int qi = QI_DT3;
+};
+
 template<>
 struct ggml_cuda_type_traits<GGML_TYPE_Q4_0> {
     static constexpr int qk = QK4_0;

ggml/src/ggml-cuda/convert.cu

@@ -461,6 +461,8 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) {
             return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
         case GGML_TYPE_Q2_0:
             return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
+        case GGML_TYPE_DT3:
+            return dequantize_block_cont_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
         case GGML_TYPE_Q4_0:
             return dequantize_row_q4_0_cuda;
         case GGML_TYPE_Q4_1:
@@ -518,6 +520,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
             return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
         case GGML_TYPE_Q2_0:
             return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
+        case GGML_TYPE_DT3:
+            return dequantize_block_cont_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
         case GGML_TYPE_Q4_0:
             return dequantize_row_q4_0_cuda;
         case GGML_TYPE_Q4_1:
@@ -578,6 +582,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
             return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
         case GGML_TYPE_Q2_0:
             return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
+        case GGML_TYPE_DT3:
+            return dequantize_block_cont_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
         case GGML_TYPE_Q4_0:
             return dequantize_row_q4_0_cuda;
         case GGML_TYPE_Q4_1:
@@ -637,6 +643,8 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) {
             return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
         case GGML_TYPE_Q2_0:
             return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
+        case GGML_TYPE_DT3:
+            return dequantize_block_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
         case GGML_TYPE_Q4_0:
             return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
         case GGML_TYPE_Q4_1:
@@ -662,6 +670,8 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) {
             return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
         case GGML_TYPE_Q2_0:
             return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
+        case GGML_TYPE_DT3:
+            return dequantize_block_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
         case GGML_TYPE_Q4_0:
             return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
         case GGML_TYPE_Q4_1:
@@ -687,6 +697,8 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) {
             return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
         case GGML_TYPE_Q2_0:
             return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
+        case GGML_TYPE_DT3:
+            return dequantize_block_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
         case GGML_TYPE_Q4_0:
             return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
         case GGML_TYPE_Q4_1:

ggml/src/ggml-cuda/dequantize.cuh

@@ -43,6 +43,19 @@ static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const in
     v.y = (c1 - 1) * d;
 }
 
+static __device__ __forceinline__ void dequantize_dt3(const void * vx, const int64_t ib, const int iqs, float2 & v){
+    const block_dt3 * x = (const block_dt3 *) vx;
+
+    // DT3: two packed ternary planes with one scale each, w = d1*t1 + d2*t2
+    const float d1 = x[ib].d[0];
+    const float d2 = x[ib].d[1];
+
+    v.x = d1*ggml_cuda_dt3_get_trit(x[ib].qs[0], x[ib].qh[0], iqs + 0) +
+          d2*ggml_cuda_dt3_get_trit(x[ib].qs[1], x[ib].qh[1], iqs + 0);
+    v.y = d1*ggml_cuda_dt3_get_trit(x[ib].qs[0], x[ib].qh[0], iqs + 1) +
+          d2*ggml_cuda_dt3_get_trit(x[ib].qs[1], x[ib].qh[1], iqs + 1);
+}
+
 static __device__ __forceinline__ void dequantize_q4_0(const void * vx, const int64_t ib, const int iqs, float2 & v){
     const block_q4_0 * x = (const block_q4_0 *) vx;
 

ggml/src/ggml-cuda/getrows.cu

@@ -324,6 +324,10 @@ static void ggml_cuda_get_rows_switch_src0_type(
             get_rows_cuda_q<QK2_0, QR2_0, dequantize_q2_0>(src0_d, src1_d, dst_d,
                 ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
             break;
+        case GGML_TYPE_DT3:
+            get_rows_cuda_q<QK_DT3, QR_DT3, dequantize_dt3>(src0_d, src1_d, dst_d,
+                ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
+            break;
         case GGML_TYPE_Q4_0:
             get_rows_cuda_q<QK4_0, QR4_0, dequantize_q4_0>(src0_d, src1_d, dst_d,
                 ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -4908,6 +4908,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
                     case GGML_TYPE_F16:
                     case GGML_TYPE_Q1_0:
                     case GGML_TYPE_Q2_0:
+                    case GGML_TYPE_DT3:
                     case GGML_TYPE_Q4_0:
                     case GGML_TYPE_Q4_1:
                     case GGML_TYPE_Q5_0:
@@ -4947,6 +4948,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
                     case GGML_TYPE_I32:
                     case GGML_TYPE_Q1_0:
                     case GGML_TYPE_Q2_0:
+                    case GGML_TYPE_DT3:
                     case GGML_TYPE_Q4_0:
                     case GGML_TYPE_Q4_1:
                     case GGML_TYPE_Q5_0:

ggml/src/ggml-cuda/mmvq.cu

@@ -11,6 +11,7 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type)
     switch (type) {
         case GGML_TYPE_Q1_0:    return vec_dot_q1_0_q8_1;
         case GGML_TYPE_Q2_0:    return vec_dot_q2_0_q8_1;
+        case GGML_TYPE_DT3:     return vec_dot_dt3_q8_1;
         case GGML_TYPE_Q4_0:    return vec_dot_q4_0_q8_1;
         case GGML_TYPE_Q4_1:    return vec_dot_q4_1_q8_1;
         case GGML_TYPE_Q5_0:    return vec_dot_q5_0_q8_1;
@@ -40,6 +41,7 @@ static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) {
     switch (type) {
         case GGML_TYPE_Q1_0:    return VDR_Q1_0_Q8_1_MMVQ;
         case GGML_TYPE_Q2_0:    return VDR_Q2_0_Q8_1_MMVQ;
+        case GGML_TYPE_DT3:     return VDR_DT3_Q8_1_MMVQ;
         case GGML_TYPE_Q4_0:    return VDR_Q4_0_Q8_1_MMVQ;
         case GGML_TYPE_Q4_1:    return VDR_Q4_1_Q8_1_MMVQ;
         case GGML_TYPE_Q5_0:    return VDR_Q5_0_Q8_1_MMVQ;
@@ -1018,6 +1020,12 @@ static void mul_mat_vec_q_switch_type(
                  nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst,
                  nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream);
             break;
+        case GGML_TYPE_DT3:
+            mul_mat_vec_q_switch_ncols_dst<GGML_TYPE_DT3>
+                (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst,
+                 nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst,
+                 nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream);
+            break;
         case GGML_TYPE_Q4_0:
             mul_mat_vec_q_switch_ncols_dst<GGML_TYPE_Q4_0>
                 (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst,

ggml/src/ggml-cuda/vecdotq.cuh

@@ -112,6 +112,8 @@ static __device__ __forceinline__ uint32_t unpack_ksigns(const uint8_t v) {
 #define VDR_Q2_0_Q8_1_MMVQ 1  // Process one 32-element chunk at a time for parallelism
 #define VDR_Q2_0_Q8_1_MMQ  2  // Q2_0 group 64: 128 bits (4 ints) per block, 2 32-element chunks
 
+#define VDR_DT3_Q8_1_MMVQ 4  // DT3: one call processes a whole 128-element block, i.e. all 4 q8_1 chunks it spans
+
 #define VDR_Q4_0_Q8_1_MMVQ 2
 #define VDR_Q4_0_Q8_1_MMQ  4
 
@@ -763,6 +765,52 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1(
     return d2 * d8 * sumi;
 }
 
+static __device__ __forceinline__ float vec_dot_dt3_q8_1(
+    const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
+
+    const block_dt3 * bq_dt3 = (const block_dt3 *) vbq + kbx;
+
+    // DT3: 128 elements as two ternary planes with one scale each, w = d1*t1 + d2*t2.
+    // One call processes the whole block (VDR_DT3_Q8_1_MMVQ == 4), so iqs is always 0
+    // and bq8_1 points to the 4 q8_1 blocks the DT3 block spans. All element indices
+    // below are compile-time constants, so the decode folds into shifts and masks.
+    GGML_UNUSED(iqs);
+
+    int sumi1[4] = {0, 0, 0, 0};
+    int sumi2[4] = {0, 0, 0, 0};
+
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+#pragma unroll
+        for (int k = 0; k < 8; ++k) {
+            const int u = get_int_b4(bq8_1[j].qs, k);
+
+            int v1 = 0;
+            int v2 = 0;
+#pragma unroll
+            for (int l = 0; l < 4; ++l) {
+                const int i = 32*j + 4*k + l;
+                v1 |= (ggml_cuda_dt3_get_trit(bq_dt3->qs[0], bq_dt3->qh[0], i) & 0xFF) << (8*l);
+                v2 |= (ggml_cuda_dt3_get_trit(bq_dt3->qs[1], bq_dt3->qh[1], i) & 0xFF) << (8*l);
+            }
+
+            sumi1[j] = ggml_cuda_dp4a(v1, u, sumi1[j]);
+            sumi2[j] = ggml_cuda_dp4a(v2, u, sumi2[j]);
+        }
+    }
+
+    const float d1 = bq_dt3->d[0];
+    const float d2 = bq_dt3->d[1];
+
+    float sumf = 0.0f;
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        const float d8 = __low2float(bq8_1[j].ds);
+        sumf += d8 * (d1*sumi1[j] + d2*sumi2[j]);
+    }
+    return sumf;
+}
+
 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) {
 

tests/CMakeLists.txt

@@ -290,6 +290,7 @@ if (NOT GGML_BACKEND_DL)
     llama_build_and_test(test-barrier.cpp)
     llama_build_and_test(test-quantize-fns.cpp)
     llama_build_and_test(test-dt3.cpp)
+    llama_build_and_test(test-dt3-gpu.cpp)
     llama_build_and_test(test-quantize-perf.cpp)
     llama_build_and_test(test-rope.cpp)
     llama_build_and_test(test-col2im-1d.cpp)

tests/test-dt3-gpu.cpp

@@ -0,0 +1,372 @@
+// GPU vs CPU parity tests for the DT3 dual-plane ternary format
+//
+// The CPU path (dequantize_row_dt3) is the validated reference. This test
+// checks the GPU backend against it in two steps:
+//
+//   1. dequantization: GET_ROWS on the GPU must reproduce the CPU reference
+//      bit by bit — same fp16 scales, exact products by {-1, 0, +1}, one
+//      float rounding per element on both sides.
+//   2. matrix multiplication: MUL_MAT with a small number of destination
+//      columns takes the MMVQ path (vec_dot_dt3_q8_1). The activations are
+//      chosen so that their q8_1 quantization is exact (integer values with
+//      amax 127 in every 32-element chunk), which makes a double precision
+//      reference computed from the dequantized weights valid to float
+//      rounding of the accumulation. One case is also checked against a
+//      manual sum over trits stored by the test, with non-trivial qh trits.
+//
+// The directed blocks exercise the three packing regions, the 79/80 and
+// 119/120 region boundaries, and negative scales. The random blocks use raw
+// random bytes: every byte value 0..255 must decode identically on both
+// sides, including values >= 243 that never come out of the packer.
+//
+// Without a GPU backend the test is skipped and succeeds.
+
+#include "ggml.h"
+#include "ggml-alloc.h"
+#include "ggml-backend.h"
+#include "ggml-cpu.h"
+
+#undef NDEBUG
+#include <assert.h>
+#include <math.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <string.h>
+#include <vector>
+
+constexpr int    QK_DT3         = 128;
+constexpr size_t DT3_QS_BYTES   = 24; // per plane
+constexpr size_t DT3_QH_BYTES   = 2;  // per plane
+constexpr size_t DT3_BLOCK_SIZE = 2*DT3_QS_BYTES + 2*DT3_QH_BYTES + 2*sizeof(uint16_t);
+
+// byte offsets inside a block (spec: qs[2][24] | qh[2][2] | d[2])
+constexpr size_t OFF_QS = 0;
+constexpr size_t OFF_QH = 2*DT3_QS_BYTES;
+constexpr size_t OFF_D  = 2*DT3_QS_BYTES + 2*DT3_QH_BYTES;
+
+// independent packer, written from the format specification (same as in
+// test-dt3.cpp): element i of a plane goes to
+//   region A: qs[m],    m in [0,16), digit n: elements m + n*16     (0..79)
+//   region B: qs[16+m], m in [0,8),  digit n: elements 80 + m + n*8 (80..119)
+//   region C: qh[j],    j in [0,2),  digit n: elements 120 + j + n*2 (120..127)
+static void ref_pack_plane(const int8_t * t, uint8_t * qs, uint8_t * qh) {
+    for (int m = 0; m < 16; ++m) {
+        uint32_t q = 0;
+        for (int n = 0; n < 5; ++n) {
+            q = q*3 + (uint32_t)(t[m + n*16] + 1);
+        }
+        qs[m] = (uint8_t)((q*256 + 242)/243);
+    }
+    for (int m = 0; m < 8; ++m) {
+        uint32_t q = 0;
+        for (int n = 0; n < 5; ++n) {
+            q = q*3 + (uint32_t)(t[80 + m + n*8] + 1);
+        }
+        qs[16 + m] = (uint8_t)((q*256 + 242)/243);
+    }
+    for (int j = 0; j < 2; ++j) {
+        uint32_t q = 0;
+        for (int n = 0; n < 4; ++n) {
+            q = q*3 + (uint32_t)(t[120 + j + n*2] + 1);
+        }
+        q *= 3; // shift the first value to the most significant trit
+        qh[j] = (uint8_t)((q*256 + 242)/243);
+    }
+}
+
+static void ref_pack_block(const int8_t * t1, float d1, const int8_t * t2, float d2, uint8_t * block) {
+    ref_pack_plane(t1, block + OFF_QS,                block + OFF_QH);
+    ref_pack_plane(t2, block + OFF_QS + DT3_QS_BYTES, block + OFF_QH + DT3_QH_BYTES);
+    const uint16_t h1 = ggml_fp32_to_fp16(d1);
+    const uint16_t h2 = ggml_fp32_to_fp16(d2);
+    memcpy(block + OFF_D,     &h1, sizeof(h1));
+    memcpy(block + OFF_D + 2, &h2, sizeof(h2));
+}
+
+// deterministic PRNG so failures are reproducible
+static uint32_t rng_state = 0x2b992ddf;
+static uint32_t rng_next(void) {
+    rng_state ^= rng_state << 13;
+    rng_state ^= rng_state >> 17;
+    rng_state ^= rng_state << 5;
+    return rng_state;
+}
+static int8_t rng_trit(void) {
+    return (int8_t)(rng_next() % 3) - 1;
+}
+
+constexpr int NROWS     = 16;
+constexpr int NCOLS     = 896; // 7 blocks per row; deliberately not a multiple of 256
+constexpr int NBLOCKS   = NROWS*NCOLS/QK_DT3;
+constexpr int ROW0_NB   = NCOLS/QK_DT3;
+
+// trits and scales of row 0, kept for the manual MUL_MAT reference
+static int8_t row0_t1[ROW0_NB][QK_DT3];
+static int8_t row0_t2[ROW0_NB][QK_DT3];
+static float  row0_d1[ROW0_NB];
+static float  row0_d2[ROW0_NB];
+
+static void build_dt3_data(std::vector<uint8_t> & data) {
+    data.resize((size_t)NBLOCKS*DT3_BLOCK_SIZE);
+
+    // row 0: known trits with non-trivial qh region and mixed-sign scales
+    for (int j = 0; j < ROW0_NB; ++j) {
+        for (int i = 0; i < QK_DT3; ++i) {
+            row0_t1[j][i] = rng_trit();
+            row0_t2[j][i] = rng_trit();
+        }
+        // make sure the qh-packed elements are not all zero
+        row0_t1[j][127] = -1;
+        row0_t2[j][120] = +1;
+
+        row0_d1[j] = j % 2 == 0 ?  1.5f  : -0.75f;  // exact in fp16
+        row0_d2[j] = j % 2 == 0 ? -0.625f:  0.375f; // exact in fp16
+
+        ref_pack_block(row0_t1[j], row0_d1[j], row0_t2[j], row0_d2[j], data.data() + (size_t)j*DT3_BLOCK_SIZE);
+    }
+
+    // directed single-trit blocks at the region boundaries, negative d2
+    const int special_pos[] = {0, 15, 16, 79, 80, 87, 88, 119, 120, 121, 126, 127};
+    const int n_special     = (int)(sizeof(special_pos)/sizeof(special_pos[0]));
+    for (int c = 0; c < n_special; ++c) {
+        int8_t t1[QK_DT3] = {0};
+        int8_t t2[QK_DT3] = {0};
+        t1[special_pos[c]] = +1;
+        t2[special_pos[c]] = -1;
+        ref_pack_block(t1, 1.0f, t2, -0.25f, data.data() + (size_t)(ROW0_NB + c)*DT3_BLOCK_SIZE);
+    }
+
+    // the rest: raw random bytes (any byte value is decodable) and random
+    // small scales, some negative
+    for (int b = ROW0_NB + n_special; b < NBLOCKS; ++b) {
+        uint8_t * block = data.data() + (size_t)b*DT3_BLOCK_SIZE;
+        for (size_t k = 0; k < OFF_D; ++k) {
+            block[k] = (uint8_t)(rng_next() & 0xFF);
+        }
+        const uint16_t h1 = ggml_fp32_to_fp16(((int)(rng_next() % 2001) - 1000)/500.0f);
+        const uint16_t h2 = ggml_fp32_to_fp16(((int)(rng_next() % 2001) - 1000)/500.0f);
+        memcpy(block + OFF_D,     &h1, sizeof(h1));
+        memcpy(block + OFF_D + 2, &h2, sizeof(h2));
+    }
+}
+
+// run a single-output graph on the backend and read the result back
+static void compute_graph(ggml_backend_t backend, ggml_context * ctx, ggml_tensor * out, float * result) {
+    ggml_cgraph * gf = ggml_new_graph(ctx);
+    ggml_build_forward_expand(gf, out);
+
+    ggml_gallocr_t galloc = ggml_gallocr_new(ggml_backend_get_default_buffer_type(backend));
+    const bool ok = ggml_gallocr_alloc_graph(galloc, gf);
+    GGML_ASSERT(ok);
+
+    const ggml_status status = ggml_backend_graph_compute(backend, gf);
+    GGML_ASSERT(status == GGML_STATUS_SUCCESS);
+
+    ggml_backend_tensor_get(out, result, 0, ggml_nbytes(out));
+    ggml_gallocr_free(galloc);
+}
+
+// GET_ROWS over all rows on the GPU vs the CPU reference dequantization
+static int test_dequant(ggml_backend_t backend, const std::vector<uint8_t> & data, const std::vector<float> & ref) {
+    ggml_init_params params = {
+        /*.mem_size   =*/ ggml_tensor_overhead()*8 + ggml_graph_overhead(),
+        /*.mem_buffer =*/ nullptr,
+        /*.no_alloc   =*/ true,
+    };
+    ggml_context * ctx = ggml_init(params);
+
+    ggml_tensor * a    = ggml_new_tensor_2d(ctx, GGML_TYPE_DT3, NCOLS, NROWS);
+    ggml_tensor * rows = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, NROWS);
+    ggml_tensor * out  = ggml_get_rows(ctx, a, rows);
+
+    if (!ggml_backend_supports_op(backend, out)) {
+        printf("FAILED: backend does not support GET_ROWS on DT3\n");
+        ggml_free(ctx);
+        return 1;
+    }
+
+    ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, backend);
+    GGML_ASSERT(buf != nullptr);
+
+    std::vector<int32_t> row_idx(NROWS);
+    for (int r = 0; r < NROWS; ++r) {
+        row_idx[r] = r;
+    }
+    ggml_backend_tensor_set(a,    data.data(),    0, data.size());
+    ggml_backend_tensor_set(rows, row_idx.data(), 0, NROWS*sizeof(int32_t));
+
+    std::vector<float> gpu((size_t)NROWS*NCOLS);
+    compute_graph(backend, ctx, out, gpu.data());
+
+    int    num_failed = 0;
+    double max_diff   = 0.0;
+    for (size_t i = 0; i < gpu.size(); ++i) {
+        const double diff = fabs((double)gpu[i] - (double)ref[i]);
+        max_diff = diff > max_diff ? diff : max_diff;
+        if (gpu[i] != ref[i]) {
+            if (num_failed < 8) {
+                printf("FAILED: dequant mismatch at block %zu elem %zu: gpu %.9g, cpu %.9g\n",
+                       i/QK_DT3, i%QK_DT3, gpu[i], ref[i]);
+            }
+            num_failed++;
+        }
+    }
+    printf("%s: dequant GPU vs CPU on %d blocks: %d mismatches, max |diff| = %g\n",
+           num_failed == 0 ? "OK" : "FAILED", NBLOCKS, num_failed, max_diff);
+
+    ggml_backend_buffer_free(buf);
+    ggml_free(ctx);
+    return num_failed == 0 ? 0 : 1;
+}
+
+// MUL_MAT on the GPU vs a double precision reference from the CPU-dequantized
+// weights. n_cols_dst <= 8 goes through MMVQ; the activations are integers
+// with amax 127 in every 32-element chunk, so their q8_1 quantization is
+// exact and the reference is valid to float accumulation rounding.
+static int test_mul_mat(ggml_backend_t backend, const std::vector<uint8_t> & data, const std::vector<float> & ref_w) {
+    int num_failed = 0;
+
+    const int ncols_dst[] = {1, 2, 5, 8, 16};
+
+    std::vector<float> y((size_t)NCOLS*16);
+    for (size_t i = 0; i < y.size(); ++i) {
+        y[i] = i % 32 == 0 ? 127.0f : (float)((int)(rng_next() % 255) - 127);
+    }
+
+    std::vector<std::vector<float>> results;
+
+    for (int c = 0; c < (int)(sizeof(ncols_dst)/sizeof(ncols_dst[0])); ++c) {
+        const int n = ncols_dst[c];
+
+        ggml_init_params params = {
+            /*.mem_size   =*/ ggml_tensor_overhead()*8 + ggml_graph_overhead(),
+            /*.mem_buffer =*/ nullptr,
+            /*.no_alloc   =*/ true,
+        };
+        ggml_context * ctx = ggml_init(params);
+
+        ggml_tensor * a   = ggml_new_tensor_2d(ctx, GGML_TYPE_DT3, NCOLS, NROWS);
+        ggml_tensor * b   = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, NCOLS, n);
+        ggml_tensor * out = ggml_mul_mat(ctx, a, b);
+
+        if (!ggml_backend_supports_op(backend, out)) {
+            printf("FAILED: backend does not support MUL_MAT on DT3\n");
+            ggml_free(ctx);
+            return 1;
+        }
+
+        ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, backend);
+        GGML_ASSERT(buf != nullptr);
+
+        ggml_backend_tensor_set(a, data.data(), 0, data.size());
+        ggml_backend_tensor_set(b, y.data(),    0, (size_t)NCOLS*n*sizeof(float));
+
+        std::vector<float> gpu((size_t)NROWS*n);
+        compute_graph(backend, ctx, out, gpu.data());
+        results.push_back(gpu);
+
+        // reference in double from the dequantized weights
+        double max_rel = 0.0;
+        for (int j = 0; j < n; ++j) {
+            for (int r = 0; r < NROWS; ++r) {
+                double sum = 0.0;
+                for (int k = 0; k < NCOLS; ++k) {
+                    sum += (double)ref_w[(size_t)r*NCOLS + k] * (double)y[(size_t)j*NCOLS + k];
+                }
+                const double rel = fabs((double)gpu[(size_t)j*NROWS + r] - sum) / (fabs(sum) > 1.0 ? fabs(sum) : 1.0);
+                max_rel = rel > max_rel ? rel : max_rel;
+            }
+        }
+        // n <= 8 is the MMVQ path with exact integer dot products; larger n
+        // falls back to dequantization + GEMM, which may run in fp16
+        const double tol = n <= 8 ? 1e-5 : 5e-3;
+        printf("%s: mul_mat GPU vs reference, ncols_dst = %2d (%s): max rel err = %g\n",
+               max_rel <= tol ? "OK" : "FAILED", n, n <= 8 ? "MMVQ" : "GEMM", max_rel);
+        if (max_rel > tol) {
+            num_failed++;
+        }
+
+        ggml_backend_buffer_free(buf);
+        ggml_free(ctx);
+    }
+
+    // MMVQ vs the dequantization-based path: first 8 columns of the GEMM run
+    // must match the ncols_dst = 8 MMVQ run
+    {
+        const std::vector<float> & mmvq = results[3]; // n = 8
+        const std::vector<float> & gemm = results[4]; // n = 16
+        double max_rel = 0.0;
+        for (int j = 0; j < 8; ++j) {
+            for (int r = 0; r < NROWS; ++r) {
+                const double v0 = mmvq[(size_t)j*NROWS + r];
+                const double v1 = gemm[(size_t)j*NROWS + r];
+                const double rel = fabs(v0 - v1) / (fabs(v0) > 1.0 ? fabs(v0) : 1.0);
+                max_rel = rel > max_rel ? rel : max_rel;
+            }
+        }
+        printf("%s: MMVQ vs GEMM path on shared columns: max rel err = %g\n",
+               max_rel <= 5e-3 ? "OK" : "FAILED", max_rel);
+        if (max_rel > 5e-3) {
+            num_failed++;
+        }
+    }
+
+    // manual sum over the trits stored by the test for row 0, column 0 —
+    // computed from the trits themselves, not from any dequantization, with
+    // non-trivial qh trits in every block of the row
+    {
+        double sum = 0.0;
+        for (int j = 0; j < ROW0_NB; ++j) {
+            for (int i = 0; i < QK_DT3; ++i) {
+                sum += (double)y[(size_t)j*QK_DT3 + i] *
+                       ((double)row0_d1[j]*row0_t1[j][i] + (double)row0_d2[j]*row0_t2[j][i]);
+            }
+        }
+        const double got = results[0][0]; // ncols_dst = 1, row 0
+        const double rel = fabs(got - sum) / (fabs(sum) > 1.0 ? fabs(sum) : 1.0);
+        printf("%s: MMVQ vs manual trit sum (row 0, col 0): gpu %.9g, manual %.9g, rel err = %g\n",
+               rel <= 1e-5 ? "OK" : "FAILED", got, sum, rel);
+        if (rel > 1e-5) {
+            num_failed++;
+        }
+    }
+
+    return num_failed;
+}
+
+int main(void) {
+    ggml_backend_t backend = nullptr;
+    for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
+        ggml_backend_dev_t dev = ggml_backend_dev_get(i);
+        if (ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_GPU) {
+            backend = ggml_backend_dev_init(dev, nullptr);
+            printf("using GPU backend: %s\n", ggml_backend_dev_name(dev));
+            break;
+        }
+    }
+    if (backend == nullptr) {
+        printf("no GPU backend available, skipping\n");
+        return 0;
+    }
+
+    std::vector<uint8_t> data;
+    build_dt3_data(data);
+
+    // CPU reference dequantization — the validated path
+    std::vector<float> ref((size_t)NROWS*NCOLS);
+    const ggml_type_traits * qfns = ggml_get_type_traits(GGML_TYPE_DT3);
+    qfns->to_float(data.data(), ref.data(), (int64_t)NROWS*NCOLS);
+
+    int num_failed = 0;
+    num_failed += test_dequant(backend, data, ref);
+    num_failed += test_mul_mat(backend, data, ref);
+
+    ggml_backend_free(backend);
+
+    if (num_failed > 0) {
+        printf("%d tests FAILED\n", num_failed);
+        return 1;
+    }
+    printf("all tests OK\n");
+    return 0;
+}