tests : skip test-dt3-gpu on backends that do not implement DT3

The test picked the first GPU device it found and treated an unsupported
op as a failure. Vulkan and SYCL answer supports_op == false for DT3,
which is the right answer for them and not a bug to report, so the test
went red on machines that were behaving correctly. Metal is worse: it
answers true for almost any type but has no DT3 shader, so the run died
in pipeline compilation halfway through.

Pick the backend by name instead — CUDA and HIP (which reports itself as
ROCm) are the only ones implementing DT3 — and skip everything else. A
supports_op failure on those two is still a real failure.

Also document that the n <= 8 gating mirrors MMVQ_MAX_BATCH_SIZE by hand
and goes stale silently if the MMVQ dispatch changes.

ggml/include/ggml.h

@@ -430,7 +430,8 @@ extern "C" {
         GGML_TYPE_NVFP4   = 40, // NVFP4 (4 blocks, E4M3 scale)
         GGML_TYPE_Q1_0    = 41,
         GGML_TYPE_Q2_0    = 42,
-        GGML_TYPE_COUNT   = 43,
+        GGML_TYPE_DT3     = 43, // DT3 (dual-plane ternary)
+        GGML_TYPE_COUNT   = 44,
     };
 
     // precision
@@ -475,6 +476,7 @@ extern "C" {
         GGML_FTYPE_MOSTLY_NVFP4   = 26, // except 1d tensors
         GGML_FTYPE_MOSTLY_Q1_0    = 27, // except 1d tensors
         GGML_FTYPE_MOSTLY_Q2_0    = 28, // except 1d tensors
+        GGML_FTYPE_MOSTLY_DT3     = 29, // except 1d tensors
     };
 
     // available tensor operations:

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
@@ -287,6 +290,26 @@ typedef struct {
 } block_tq2_0;
 static_assert(sizeof(block_tq2_0) == sizeof(ggml_half) + QK_K / 4, "wrong tq2_0 block size/padding");
 
+//
+// Dual-plane ternary quantization (DT3)
+//
+// w_i = d[0]*t0_i + d[1]*t1_i, with t in {-1, 0, +1}. Two ternary planes buy
+// back most of the quality a single one loses, at 3.5 bpw instead of 1.6875.
+//
+// Each plane is packed exactly like tq1_0 with every constant halved, since the
+// block is 128 elements instead of QK_K: qs 48 -> 24 bytes, qh 4 -> 2 bytes,
+// and the two qs passes go over 16 and then 8 bytes instead of 32 and 16.
+//
+#define QK_DT3 128
+
+// 3.5 bpw
+typedef struct {
+    uint8_t qs[2][(QK_DT3 - 4 * QK_DT3 / 64) / 5]; // 2 planes, 5 elements per byte (3^5 = 243 < 256)
+    uint8_t qh[2][QK_DT3 / 64];                    // 2 planes, 4 elements per byte
+    ggml_half d[2];                                // one scale per plane
+} block_dt3;
+static_assert(sizeof(block_dt3) == 2 * (sizeof(ggml_half) + QK_DT3 / 64 + (QK_DT3 - 4 * QK_DT3 / 64) / 5), "wrong dt3 block size/padding");
+
 //
 // Super-block quantization structures
 //

ggml/src/ggml-cpu/arch-fallback.h

@@ -20,6 +20,7 @@
 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0
 #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K
 #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K
+#define ggml_vec_dot_dt3_q8_0_generic ggml_vec_dot_dt3_q8_0
 #define ggml_vec_dot_q2_K_q8_K_generic ggml_vec_dot_q2_K_q8_K
 #define ggml_vec_dot_q3_K_q8_K_generic ggml_vec_dot_q3_K_q8_K
 #define ggml_vec_dot_q4_K_q8_K_generic ggml_vec_dot_q4_K_q8_K
@@ -72,6 +73,8 @@
 #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0
 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0
 #elif defined(__aarch64__) || defined(__arm__) || defined(_M_ARM) || defined(_M_ARM64)
+// quants.c
+#define ggml_vec_dot_dt3_q8_0_generic ggml_vec_dot_dt3_q8_0
 // repack.cpp
 #define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4
 #define ggml_quantize_mat_q8_K_4x8_generic ggml_quantize_mat_q8_K_4x8
@@ -118,6 +121,7 @@
 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0
 #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K
 #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K
+#define ggml_vec_dot_dt3_q8_0_generic ggml_vec_dot_dt3_q8_0
 #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K
 // repack.cpp
 #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4
@@ -161,6 +165,7 @@
 #define quantize_row_q8_K_generic quantize_row_q8_K
 #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K
 #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K
+#define ggml_vec_dot_dt3_q8_0_generic ggml_vec_dot_dt3_q8_0
 #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K
 #define ggml_vec_dot_mxfp4_q8_0_generic ggml_vec_dot_mxfp4_q8_0
 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0
@@ -207,6 +212,7 @@
 // quants.c
 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0
 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0
+#define ggml_vec_dot_dt3_q8_0_generic ggml_vec_dot_dt3_q8_0
 // repack.cpp
 #define ggml_quantize_mat_q8_0_4x1_generic ggml_quantize_mat_q8_0_4x1
 #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4
@@ -251,6 +257,7 @@
 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0
 #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K
 #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K
+#define ggml_vec_dot_dt3_q8_0_generic ggml_vec_dot_dt3_q8_0
 #define ggml_vec_dot_q2_K_q8_K_generic ggml_vec_dot_q2_K_q8_K
 #define ggml_vec_dot_iq2_xxs_q8_K_generic ggml_vec_dot_iq2_xxs_q8_K
 #define ggml_vec_dot_iq2_xs_q8_K_generic ggml_vec_dot_iq2_xs_q8_K
@@ -300,6 +307,7 @@
 // quants.c
 #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K
 #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K
+#define ggml_vec_dot_dt3_q8_0_generic ggml_vec_dot_dt3_q8_0
 #define ggml_vec_dot_iq2_xxs_q8_K_generic ggml_vec_dot_iq2_xxs_q8_K
 #define ggml_vec_dot_iq2_xs_q8_K_generic ggml_vec_dot_iq2_xs_q8_K
 #define ggml_vec_dot_iq2_s_q8_K_generic ggml_vec_dot_iq2_s_q8_K

ggml/src/ggml-cpu/arch/x86/quants.c

@@ -1571,6 +1571,179 @@ void ggml_vec_dot_tq2_0_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const vo
 #endif
 }
 
+#if defined(__AVX512F__) && defined(__AVX512BW__) && defined(__AVX512VBMI__) && defined(__AVX512VNNI__)
+// multiply bytes by 3 with wraparound (there is no 8-bit SIMD multiply)
+static inline __m512i dt3_mul3_epi8(const __m512i v) {
+    return _mm512_add_epi8(v, _mm512_add_epi8(v, v));
+}
+
+// bring the top base-3 digit of each byte down to xi = ((uint8_t) q * 3) >> 8,
+// in {0, 1, 2}, with the same avg trick as ggml_vec_dot_tq1_0_q8_K
+static inline __m512i dt3_decode_epi8(__m512i q) {
+    // cancel the +1 from avg so that it behaves like a halving add
+    q = _mm512_subs_epu8(q, _mm512_set1_epi8(1));
+    // multiply by 3 and get the top 2 bits
+    q = _mm512_avg_epu8(q, _mm512_avg_epu8(q, _mm512_setzero_si512()));
+    return _mm512_and_si512(_mm512_srli_epi16(q, 6), _mm512_set1_epi8(3));
+}
+#endif
+
+void ggml_vec_dot_dt3_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
+    assert(n % QK_DT3 == 0);
+    assert(nrc == 1);
+    UNUSED(nrc);
+    UNUSED(bx);
+    UNUSED(by);
+    UNUSED(bs);
+
+    const block_dt3  * GGML_RESTRICT x = vx;
+    const block_q8_0 * GGML_RESTRICT y = vy;
+
+    const int nb = n / QK_DT3;
+
+#if defined(__AVX512F__) && defined(__AVX512BW__) && defined(__AVX512VBMI__) && defined(__AVX512VNNI__)
+    // Source byte of each of the 128 elements of a plane, in element order, as
+    // offsets into the 56-byte block for plane 0 (see unpack_plane_dt3):
+    //   element m + n*16      (0..79)    is digit n of qs[m],      m in [0,16)
+    //   element 80 + m + n*8  (80..119)  is digit n of qs[16 + m], m in [0,8)
+    //   element 120 + j + n*2 (120..127) is digit n of qh[j],      j in [0,2)
+    // The low vector covers elements 0..63 and the high vector 64..127, so
+    // that each aligns with two full q8_0 blocks of the other operand.
+    static const uint8_t kidx_lo[64] = {
+         0,  1,  2,  3,  4,  5,  6,  7,  8,  9, 10, 11, 12, 13, 14, 15, // digit 0 of qs[0..15]
+         0,  1,  2,  3,  4,  5,  6,  7,  8,  9, 10, 11, 12, 13, 14, 15, // digit 1
+         0,  1,  2,  3,  4,  5,  6,  7,  8,  9, 10, 11, 12, 13, 14, 15, // digit 2
+         0,  1,  2,  3,  4,  5,  6,  7,  8,  9, 10, 11, 12, 13, 14, 15, // digit 3
+    };
+    static const uint8_t kidx_hi[64] = {
+         0,  1,  2,  3,  4,  5,  6,  7,  8,  9, 10, 11, 12, 13, 14, 15, // digit 4 of qs[0..15]
+        16, 17, 18, 19, 20, 21, 22, 23,                                 // digit 0 of qs[16..23]
+        16, 17, 18, 19, 20, 21, 22, 23,                                 // digit 1
+        16, 17, 18, 19, 20, 21, 22, 23,                                 // digit 2
+        16, 17, 18, 19, 20, 21, 22, 23,                                 // digit 3
+        16, 17, 18, 19, 20, 21, 22, 23,                                 // digit 4
+        48, 49, 48, 49, 48, 49, 48, 49,                                 // digits 0,0,1,1,2,2,3,3 of qh[0],qh[1]
+    };
+    // plane 1 offsets: qs starts 24 bytes later, qh starts 2 bytes later
+    static const uint8_t koff_hi[64] = {
+        24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,
+        24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,
+        24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24, 24,
+        24, 24, 24, 24, 24, 24, 24, 24,  2,  2,  2,  2,  2,  2,  2,  2,
+    };
+
+    // digit -> multiplier blend masks, one bit per byte lane (digit n of a
+    // byte is extracted by multiplying by 3^n and reading the top 2 bits;
+    // the qh lanes of the high vector never see 3^4, which would read the
+    // padding 5th digit of the qh bytes)
+    const __mmask64 mask_lo_3  = (__mmask64) 0x00000000FFFF0000; // lanes 16..31
+    const __mmask64 mask_lo_9  = (__mmask64) 0x0000FFFF00000000; // lanes 32..47
+    const __mmask64 mask_lo_27 = (__mmask64) 0xFFFF000000000000; // lanes 48..63
+    const __mmask64 mask_hi_1  = (__mmask64) 0x0300000000FF0000; // lanes 16..23, 56, 57
+    const __mmask64 mask_hi_3  = (__mmask64) 0x0C000000FF000000; // lanes 24..31, 58, 59
+    const __mmask64 mask_hi_9  = (__mmask64) 0x300000FF00000000; // lanes 32..39, 60, 61
+    const __mmask64 mask_hi_27 = (__mmask64) 0xC000FF0000000000; // lanes 40..47, 62, 63
+
+    const __m512i idx_lo0 = _mm512_loadu_si512(kidx_lo);
+    const __m512i idx_hi0 = _mm512_loadu_si512(kidx_hi);
+    const __m512i idx_lo1 = _mm512_add_epi8(idx_lo0, _mm512_set1_epi8(24));
+    const __m512i idx_hi1 = _mm512_add_epi8(idx_hi0, _mm512_loadu_si512(koff_hi));
+
+    const __m512i zero = _mm512_setzero_si512();
+    const __m512i ones = _mm512_set1_epi8(1);
+
+    float sumf = 0.0f;
+
+    for (int i = 0; i < nb; i++) {
+        // the block is 56 bytes; the masked load must not read past the end
+        const __m512i blk = _mm512_maskz_loadu_epi8((__mmask64) ((UINT64_C(1) << sizeof(block_dt3)) - 1), &x[i]);
+
+        // multiplying by 3^n (with wraparound) commutes with the byte
+        // permutation, so the multiply chain is computed once on the whole
+        // block and shared by both planes; masked vpermb then picks each
+        // lane's byte from the chain vector of its digit
+        const __m512i v3  = dt3_mul3_epi8(blk);
+        const __m512i v9  = dt3_mul3_epi8(v3);
+        const __m512i v27 = dt3_mul3_epi8(v9);
+        const __m512i v81 = dt3_mul3_epi8(v27);
+
+        __m512i xi_lo[2];
+        __m512i xi_hi[2];
+
+        for (int p = 0; p < 2; p++) {
+            const __m512i idx_lo = p == 0 ? idx_lo0 : idx_lo1;
+            const __m512i idx_hi = p == 0 ? idx_hi0 : idx_hi1;
+
+            __m512i q_lo = _mm512_permutexvar_epi8(idx_lo, blk);
+            q_lo = _mm512_mask_permutexvar_epi8(q_lo, mask_lo_3,  idx_lo, v3);
+            q_lo = _mm512_mask_permutexvar_epi8(q_lo, mask_lo_9,  idx_lo, v9);
+            q_lo = _mm512_mask_permutexvar_epi8(q_lo, mask_lo_27, idx_lo, v27);
+            xi_lo[p] = dt3_decode_epi8(q_lo);
+
+            __m512i q_hi = _mm512_permutexvar_epi8(idx_hi, v81);
+            q_hi = _mm512_mask_permutexvar_epi8(q_hi, mask_hi_1,  idx_hi, blk);
+            q_hi = _mm512_mask_permutexvar_epi8(q_hi, mask_hi_3,  idx_hi, v3);
+            q_hi = _mm512_mask_permutexvar_epi8(q_hi, mask_hi_9,  idx_hi, v9);
+            q_hi = _mm512_mask_permutexvar_epi8(q_hi, mask_hi_27, idx_hi, v27);
+            xi_hi[p] = dt3_decode_epi8(q_hi);
+        }
+
+        // one DT3 block (128 weights) maps to four q8_0 blocks (4 * 32 = 128)
+        const __m512i y_lo = _mm512_inserti64x4(_mm512_castsi256_si512(
+                                 _mm256_loadu_si256((const __m256i *) y[4*i + 0].qs)),
+                                 _mm256_loadu_si256((const __m256i *) y[4*i + 1].qs), 1);
+        const __m512i y_hi = _mm512_inserti64x4(_mm512_castsi256_si512(
+                                 _mm256_loadu_si256((const __m256i *) y[4*i + 2].qs)),
+                                 _mm256_loadu_si256((const __m256i *) y[4*i + 3].qs), 1);
+
+        // t = xi - 1, so t.y = xi.y - sum(y)
+        const __m512i sy_lo = _mm512_dpbusd_epi32(zero, ones, y_lo);
+        const __m512i sy_hi = _mm512_dpbusd_epi32(zero, ones, y_hi);
+
+        const __m512i t1_lo = _mm512_sub_epi32(_mm512_dpbusd_epi32(zero, xi_lo[0], y_lo), sy_lo);
+        const __m512i t1_hi = _mm512_sub_epi32(_mm512_dpbusd_epi32(zero, xi_hi[0], y_hi), sy_hi);
+        const __m512i t2_lo = _mm512_sub_epi32(_mm512_dpbusd_epi32(zero, xi_lo[1], y_lo), sy_lo);
+        const __m512i t2_hi = _mm512_sub_epi32(_mm512_dpbusd_epi32(zero, xi_hi[1], y_hi), sy_hi);
+
+        // reduce each q8_0 block (8 consecutive int32 lanes) to its sum with
+        // an hadd tree; integer addition order does not affect the result:
+        //   ab   = [A01 A23 B01 B23 | A45 A67 B45 B67]
+        //   abcd = [A0123 B0123 C0123 D0123 | A4567 B4567 C4567 D4567]
+        //   sv   = [sum(A) sum(B) sum(C) sum(D)]
+        const __m256i ab1  = _mm256_hadd_epi32(_mm512_castsi512_si256(t1_lo), _mm512_extracti64x4_epi64(t1_lo, 1));
+        const __m256i cd1  = _mm256_hadd_epi32(_mm512_castsi512_si256(t1_hi), _mm512_extracti64x4_epi64(t1_hi, 1));
+        const __m256i ab2  = _mm256_hadd_epi32(_mm512_castsi512_si256(t2_lo), _mm512_extracti64x4_epi64(t2_lo, 1));
+        const __m256i cd2  = _mm256_hadd_epi32(_mm512_castsi512_si256(t2_hi), _mm512_extracti64x4_epi64(t2_hi, 1));
+        const __m256i abcd1 = _mm256_hadd_epi32(ab1, cd1);
+        const __m256i abcd2 = _mm256_hadd_epi32(ab2, cd2);
+        const __m128i sv1 = _mm_add_epi32(_mm256_castsi256_si128(abcd1), _mm256_extracti128_si256(abcd1, 1));
+        const __m128i sv2 = _mm_add_epi32(_mm256_castsi256_si128(abcd2), _mm256_extracti128_si256(abcd2, 1));
+
+        int sumi1[4];
+        int sumi2[4];
+        _mm_storeu_si128((__m128i *) sumi1, sv1);
+        _mm_storeu_si128((__m128i *) sumi2, sv2);
+
+        const float d1 = GGML_CPU_FP16_TO_FP32(x[i].d[0]);
+        const float d2 = GGML_CPU_FP16_TO_FP32(x[i].d[1]);
+
+        // same accumulation order as the generic implementation
+        for (int k = 0; k < 4; k++) {
+            const float dy = GGML_CPU_FP16_TO_FP32(y[4*i + k].d);
+            sumf += dy * (d1*sumi1[k] + d2*sumi2[k]);
+        }
+    }
+
+    *s = sumf;
+
+#else
+    UNUSED(x);
+    UNUSED(y);
+    UNUSED(nb);
+    ggml_vec_dot_dt3_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc);
+#endif
+}
+
 void ggml_vec_dot_q2_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
     assert(nrc == 1);
     UNUSED(nrc);

ggml/src/ggml-cpu/ggml-cpu.c

@@ -409,6 +409,14 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = {
         .vec_dot_type             = GGML_TYPE_Q8_K,
         .nrows                    = 1,
     },
+    [GGML_TYPE_DT3] = {
+        .from_float               = quantize_row_dt3,
+        .vec_dot                  = ggml_vec_dot_dt3_q8_0,
+        // Q8_0 on purpose: DT3 has no zero-point, so the q8_K bsums are dead
+        // weight, and 32-element blocks allow any row size multiple of 128
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
     [GGML_TYPE_I32] = {
         .from_float               = (ggml_from_float_t) ggml_cpu_fp32_to_i32,
     },

ggml/src/ggml-cpu/ops.cpp

@@ -680,6 +680,7 @@ void ggml_compute_forward_add(
         case GGML_TYPE_Q6_K:
         case GGML_TYPE_TQ1_0:
         case GGML_TYPE_TQ2_0:
+        case GGML_TYPE_DT3:
         case GGML_TYPE_IQ2_XXS:
         case GGML_TYPE_IQ2_XS:
         case GGML_TYPE_IQ3_XXS:
@@ -1132,6 +1133,7 @@ void ggml_compute_forward_add1(
         case GGML_TYPE_Q6_K:
         case GGML_TYPE_TQ1_0:
         case GGML_TYPE_TQ2_0:
+        case GGML_TYPE_DT3:
         case GGML_TYPE_IQ2_XXS:
         case GGML_TYPE_IQ2_XS:
         case GGML_TYPE_IQ3_XXS:
@@ -1263,6 +1265,7 @@ void ggml_compute_forward_acc(
         case GGML_TYPE_Q6_K:
         case GGML_TYPE_TQ1_0:
         case GGML_TYPE_TQ2_0:
+        case GGML_TYPE_DT3:
         case GGML_TYPE_IQ2_XXS:
         case GGML_TYPE_IQ2_XS:
         case GGML_TYPE_IQ3_XXS:
@@ -4536,6 +4539,7 @@ void ggml_compute_forward_out_prod(
         case GGML_TYPE_Q6_K:
         case GGML_TYPE_TQ1_0:
         case GGML_TYPE_TQ2_0:
+        case GGML_TYPE_DT3:
         case GGML_TYPE_IQ2_XXS:
         case GGML_TYPE_IQ2_XS:
         case GGML_TYPE_IQ3_XXS:
@@ -4813,6 +4817,7 @@ void ggml_compute_forward_set(
         case GGML_TYPE_Q6_K:
         case GGML_TYPE_TQ1_0:
         case GGML_TYPE_TQ2_0:
+        case GGML_TYPE_DT3:
         case GGML_TYPE_IQ2_XXS:
         case GGML_TYPE_IQ2_XS:
         case GGML_TYPE_IQ3_XXS:
@@ -5038,6 +5043,7 @@ void ggml_compute_forward_get_rows(
         case GGML_TYPE_Q6_K:
         case GGML_TYPE_TQ1_0:
         case GGML_TYPE_TQ2_0:
+        case GGML_TYPE_DT3:
         case GGML_TYPE_IQ2_XXS:
         case GGML_TYPE_IQ2_XS:
         case GGML_TYPE_IQ3_XXS:
@@ -5795,6 +5801,7 @@ void ggml_compute_forward_clamp(
         case GGML_TYPE_Q6_K:
         case GGML_TYPE_TQ1_0:
         case GGML_TYPE_TQ2_0:
+        case GGML_TYPE_DT3:
         case GGML_TYPE_IQ2_XXS:
         case GGML_TYPE_IQ2_XS:
         case GGML_TYPE_IQ3_XXS:

ggml/src/ggml-cpu/quants.c

@@ -116,6 +116,12 @@ void quantize_row_tq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy,
     quantize_row_tq2_0_ref(x, y, k);
 }
 
+void quantize_row_dt3(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) {
+    assert(k % QK_DT3 == 0);
+    block_dt3 * GGML_RESTRICT y = vy;
+    quantize_row_dt3_ref(x, y, k);
+}
+
 //===================================== Q8_K ==============================================
 
 void quantize_row_q8_K_generic(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) {
@@ -562,6 +568,51 @@ void ggml_vec_dot_tq2_0_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs,
     *s = sumf;
 }
 
+void ggml_vec_dot_dt3_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
+    const int qk = QK_DT3;
+    const int nb = n / qk;
+
+    assert(n % qk == 0);
+    assert(nrc == 1);
+    UNUSED(nrc);
+    UNUSED(bx);
+    UNUSED(by);
+    UNUSED(bs);
+
+    const block_dt3  * GGML_RESTRICT x = vx;
+    const block_q8_0 * GGML_RESTRICT y = vy;
+
+    float sumf = 0.0f;
+
+    for (int i = 0; i < nb; i++) {
+        int8_t t1[QK_DT3];
+        int8_t t2[QK_DT3];
+
+        unpack_plane_dt3(x[i].qs[0], x[i].qh[0], t1);
+        unpack_plane_dt3(x[i].qs[1], x[i].qh[1], t2);
+
+        const float d1 = GGML_CPU_FP16_TO_FP32(x[i].d[0]);
+        const float d2 = GGML_CPU_FP16_TO_FP32(x[i].d[1]);
+
+        // one DT3 block (128 weights) maps to four q8_0 blocks (4 * 32 = 128)
+        for (int k = 0; k < 4; k++) {
+            const block_q8_0 * GGML_RESTRICT yb = &y[i*4 + k];
+            const float dy = GGML_CPU_FP16_TO_FP32(yb->d);
+
+            int sumi1 = 0;
+            int sumi2 = 0;
+            for (int j = 0; j < QK8_0; j++) {
+                sumi1 += t1[k*QK8_0 + j] * yb->qs[j];
+                sumi2 += t2[k*QK8_0 + j] * yb->qs[j];
+            }
+
+            sumf += dy * (d1*sumi1 + d2*sumi2);
+        }
+    }
+
+    *s = sumf;
+}
+
 void ggml_vec_dot_q2_K_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
     assert(nrc == 1);
     UNUSED(nrc);

ggml/src/ggml-cpu/quants.h

@@ -33,6 +33,7 @@ void quantize_row_q8_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in
 
 void quantize_row_tq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
 void quantize_row_tq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_dt3  (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
 
 void quantize_row_iq4_nl (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
 void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
@@ -57,6 +58,7 @@ void ggml_vec_dot_q6_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const voi
 
 void ggml_vec_dot_tq1_0_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 void ggml_vec_dot_tq2_0_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_dt3_q8_0  (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 
 void ggml_vec_dot_iq2_xxs_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 void ggml_vec_dot_iq2_xs_q8_K (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
@@ -85,6 +87,7 @@ void ggml_vec_dot_nvfp4_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs,
 
 void ggml_vec_dot_tq1_0_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 void ggml_vec_dot_tq2_0_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_dt3_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 
 void ggml_vec_dot_q2_K_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 void ggml_vec_dot_q3_K_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);

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,113 @@ 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);
+
+    // The dot product is accumulated over base-3 digits in {0, 1, 2} instead of
+    // trits in {-1, 0, +1}: sum((digit - 1)*u) == sum(digit*u) - sum(u), and sum(u)
+    // is one extra dp4a with 0x01010101 shared by both planes. This allows decoding
+    // each packed byte once with the iteration q -> (q*3) & 0xFF, whose step n
+    // exposes base-3 digit n of the byte in bits 8..9 of q*3, instead of
+    // re-multiplying the byte by a power of 3 for every one of its 5 elements.
+    // Two bytes are iterated at a time in the 16-bit lanes of one int: a lane
+    // holds q < 256, so q*3 < 768 never carries into the neighbouring lane.
+    // The subtraction sumi - sumu happens in exact integer arithmetic, so the
+    // result is bit-identical to decoding the trits one by one.
+
+    int sumi[2][4] = {{0, 0, 0, 0}, {0, 0, 0, 0}}; // per plane, per q8_1 chunk: sum(digit*u)
+    int sumu[4]    = {0, 0, 0, 0};                 // per q8_1 chunk: sum(u)
+
+#pragma unroll
+    for (int p = 0; p < 2; ++p) {
+        // qs[0..16): 4 quads of consecutive bytes, byte m holds elements m + 16*n
+#pragma unroll
+        for (int g = 0; g < 4; ++g) {
+            const int x = get_int_b4(bq_dt3->qs[p], g);
+            int qa = (x >> 0) & 0x00FF00FF; // bytes 4*g + 0 and 4*g + 2
+            int qb = (x >> 8) & 0x00FF00FF; // bytes 4*g + 1 and 4*g + 3
+#pragma unroll
+            for (int n = 0; n < 5; ++n) {
+                const int qa3 = qa*3;
+                const int qb3 = qb*3;
+                const int dig = ((qa3 >> 8) & 0x00030003) | (qb3 & 0x03000300);
+                qa = qa3 & 0x00FF00FF;
+                qb = qb3 & 0x00FF00FF;
+
+                const int i = 16*n + 4*g; // first of the 4 consecutive elements
+                const int j = i / 32;
+                const int u = get_int_b4(bq8_1[j].qs, (i % 32)/4);
+                if (p == 0) {
+                    sumu[j] = ggml_cuda_dp4a(0x01010101, u, sumu[j]);
+                }
+                sumi[p][j] = ggml_cuda_dp4a(dig, u, sumi[p][j]);
+            }
+        }
+        // qs[16..24): 2 quads, byte 16 + m holds elements 80 + m + 8*n
+#pragma unroll
+        for (int g = 0; g < 2; ++g) {
+            const int x = get_int_b4(bq_dt3->qs[p], 4 + g);
+            int qa = (x >> 0) & 0x00FF00FF;
+            int qb = (x >> 8) & 0x00FF00FF;
+#pragma unroll
+            for (int n = 0; n < 5; ++n) {
+                const int qa3 = qa*3;
+                const int qb3 = qb*3;
+                const int dig = ((qa3 >> 8) & 0x00030003) | (qb3 & 0x03000300);
+                qa = qa3 & 0x00FF00FF;
+                qb = qb3 & 0x00FF00FF;
+
+                const int i = 80 + 8*n + 4*g;
+                const int j = i / 32;
+                const int u = get_int_b4(bq8_1[j].qs, (i % 32)/4);
+                if (p == 0) {
+                    sumu[j] = ggml_cuda_dp4a(0x01010101, u, sumu[j]);
+                }
+                sumi[p][j] = ggml_cuda_dp4a(dig, u, sumi[p][j]);
+            }
+        }
+        // qh[0..2): byte b holds elements 120 + b + 2*n for n = 0..4 — only 4
+        // digits are iterated, the 5th is packing padding and must never decode
+        {
+            int q = bq_dt3->qh[p][0] | (bq_dt3->qh[p][1] << 16);
+#pragma unroll
+            for (int s = 0; s < 2; ++s) { // one q8_1 int: elements 120 + 4*s .. 123 + 4*s
+                int dig = 0;
+#pragma unroll
+                for (int n = 0; n < 2; ++n) {
+                    const int q3 = q*3;
+                    dig |= (((q3 >> 8) & 0x03) | ((q3 >> 16) & 0x0300)) << (16*n);
+                    q = q3 & 0x00FF00FF;
+                }
+                const int u = get_int_b4(bq8_1[3].qs, 6 + s);
+                if (p == 0) {
+                    sumu[3] = ggml_cuda_dp4a(0x01010101, u, sumu[3]);
+                }
+                sumi[p][3] = ggml_cuda_dp4a(dig, u, sumi[p][3]);
+            }
+        }
+    }
+
+    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*(sumi[0][j] - sumu[j]) + d2*(sumi[1][j] - sumu[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) {
 

ggml/src/ggml-quants.c

@@ -2483,6 +2483,201 @@ void dequantize_row_tq2_0(const block_tq2_0 * GGML_RESTRICT x, float * GGML_REST
     }
 }
 
+// ====================== Dual-plane ternary (de)-quantization (DT3)
+
+// packs one plane of QK_DT3 trits in {-1, 0, 1}, exactly like tq1_0 with all
+// constants halved (see block_dt3 in ggml-common.h)
+static void pack_plane_dt3(const int8_t * GGML_RESTRICT t, uint8_t * GGML_RESTRICT qs, uint8_t * GGML_RESTRICT qh) {
+    // 5 elements per byte, along 16 bytes
+    for (size_t m = 0; m < 16; ++m) {
+        uint8_t q = 0;
+        for (size_t n = 0; n < 5; ++n) {
+            const int xi = t[m + n*16] + 1; // -1, 0, 1 -> 0, 1, 2
+            q *= 3;
+            q += xi;
+        }
+        // ceiling division (243 == pow(3, 5))
+        q = ((uint16_t)q * 256 + (243 - 1)) / 243;
+        qs[m] = q;
+    }
+    // along 8 bytes
+    for (size_t m = 0; m < 8; ++m) {
+        uint8_t q = 0;
+        for (size_t n = 0; n < 5; ++n) {
+            const int xi = t[80 + m + n*8] + 1; // -1, 0, 1 -> 0, 1, 2
+            q *= 3;
+            q += xi;
+        }
+        // ceiling division (243 == pow(3, 5))
+        q = ((uint16_t)q * 256 + (243 - 1)) / 243;
+        qs[16 + m] = q;
+    }
+    // 4 elements per byte
+    for (size_t j = 0; j < 2; ++j) {
+        uint8_t q = 0;
+        for (size_t m = 0; m < 4; ++m) {
+            const int xi = t[120 + j + m*2] + 1; // -1, 0, 1 -> 0, 1, 2
+            q *= 3;
+            q += xi;
+        }
+        // shift the first value to the most significant trit
+        q *= 3;
+        // ceiling division (243 == pow(3, 5))
+        q = ((uint16_t)q * 256 + (243 - 1)) / 243;
+        qh[j] = q;
+    }
+}
+
+// NOTE: this is NOT the PTQTP coordinate-descent solver used to produce the
+// published DT3 models; that solver lives in the ternaria project and its
+// output is packed directly into this layout. This reference is a greedy
+// two-pass quantization (plane 1 by absolute max, plane 2 on the residual)
+// plus two rounds of alternating least-squares refits — good enough for the
+// type to be usable from float, but the measured quality of DT3 is only
+// obtained through the external PTQTP pipeline.
+void quantize_row_dt3_ref(const float * GGML_RESTRICT x, block_dt3 * GGML_RESTRICT y, int64_t k) {
+    assert(k % QK_DT3 == 0);
+    const int64_t nb = k / QK_DT3;
+
+    for (int64_t i = 0; i < nb; i++) {
+        int8_t t1[QK_DT3];
+        int8_t t2[QK_DT3];
+        float  d1;
+        float  d2;
+
+        // plane 1: ternary quantization by absolute max, like tq1_0
+        {
+            float amax = 0.0f;
+            for (int j = 0; j < QK_DT3; j++) {
+                amax = MAX(amax, fabsf(x[j]));
+            }
+            d1 = amax;
+            const float id = d1 ? 1.0f/d1 : 0.0f;
+            for (int j = 0; j < QK_DT3; j++) {
+                const int v = (int) lroundf(x[j] * id); // -1, 0, 1 (clamp guards against NaN input)
+                t1[j] = (int8_t) MAX(-1, MIN(1, v));
+            }
+        }
+
+        // plane 2: ternary quantization of the residual
+        {
+            float amax = 0.0f;
+            for (int j = 0; j < QK_DT3; j++) {
+                amax = MAX(amax, fabsf(x[j] - d1*t1[j]));
+            }
+            d2 = amax;
+            const float id = d2 ? 1.0f/d2 : 0.0f;
+            for (int j = 0; j < QK_DT3; j++) {
+                const int v = (int) lroundf((x[j] - d1*t1[j]) * id); // -1, 0, 1 (clamp guards against NaN input)
+                t2[j] = (int8_t) MAX(-1, MIN(1, v));
+            }
+        }
+
+        // two rounds of alternating refits: least-squares scale given the
+        // trits, then re-solve the trits given the scale, holding the other
+        // plane fixed
+        for (int it = 0; it < 2; ++it) {
+            for (int p = 0; p < 2; ++p) {
+                int8_t       * t      = p == 0 ? t1 : t2;
+                const int8_t * tother = p == 0 ? t2 : t1;
+                float        * d      = p == 0 ? &d1 : &d2;
+                const float    dother = p == 0 ? d2 : d1;
+
+                float sum_rt = 0.0f;
+                float sum_tt = 0.0f;
+                for (int j = 0; j < QK_DT3; j++) {
+                    const float r = x[j] - dother*tother[j];
+                    sum_rt += r * t[j];
+                    sum_tt += (float)(t[j] * t[j]);
+                }
+                if (sum_tt > 0.0f) {
+                    *d = sum_rt / sum_tt;
+                }
+                const float id = *d ? 1.0f/(*d) : 0.0f;
+                for (int j = 0; j < QK_DT3; j++) {
+                    const float r = x[j] - dother*tother[j];
+                    const int   v = (int) lroundf(r * id);
+                    t[j] = (int8_t) MAX(-1, MIN(1, v));
+                }
+            }
+        }
+
+        y[i].d[0] = GGML_FP32_TO_FP16(d1);
+        y[i].d[1] = GGML_FP32_TO_FP16(d2);
+
+        pack_plane_dt3(t1, y[i].qs[0], y[i].qh[0]);
+        pack_plane_dt3(t2, y[i].qs[1], y[i].qh[1]);
+
+        x += QK_DT3;
+    }
+}
+
+size_t quantize_dt3(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) {
+    if (quant_weights) {
+        // say it loudly, or an imatrix-vs-no-imatrix A/B on DT3 would come out
+        // byte-identical and someone would conclude the imatrix does nothing
+        static bool warned = false;
+        if (!warned) {
+            fprintf(stderr, "%s: WARNING: imatrix ignored: the DT3 reference quantizer does not use it; quality comes from ternaria's PTQTP\n", __func__);
+            warned = true;
+        }
+    }
+    const size_t row_size = ggml_row_size(GGML_TYPE_DT3, n_per_row);
+    quantize_row_dt3_ref(src, dst, (int64_t)nrow*n_per_row);
+    return nrow * row_size;
+}
+
+void unpack_plane_dt3(const uint8_t * GGML_RESTRICT qs, const uint8_t * GGML_RESTRICT qh, int8_t * GGML_RESTRICT t) {
+    const uint8_t pow3[6] = {1, 3, 9, 27, 81, 243};
+
+    // 5 elements per byte, along 16 bytes
+    for (size_t n = 0; n < 5; ++n) {
+        for (size_t m = 0; m < 16; ++m) {
+            const uint8_t q  = qs[m] * pow3[n]; // the product wraps around on purpose
+            const int16_t xi = ((uint16_t) q * 3) >> 8;
+            *t++ = (int8_t)(xi - 1);
+        }
+    }
+    // along 8 bytes
+    for (size_t n = 0; n < 5; ++n) {
+        for (size_t m = 0; m < 8; ++m) {
+            const uint8_t q  = qs[16 + m] * pow3[n];
+            const int16_t xi = ((uint16_t) q * 3) >> 8;
+            *t++ = (int8_t)(xi - 1);
+        }
+    }
+    // 4 elements per byte — NOT 5: each qh byte stores only 4 trits, and its
+    // 5th base-3 digit is padding from the packer's extra q *= 3 shift which
+    // always decodes to -1; reading it would inject spurious values
+    for (size_t n = 0; n < 4; ++n) {
+        for (size_t j = 0; j < 2; ++j) {
+            const uint8_t q  = qh[j] * pow3[n];
+            const int16_t xi = ((uint16_t) q * 3) >> 8;
+            *t++ = (int8_t)(xi - 1);
+        }
+    }
+}
+
+void dequantize_row_dt3(const block_dt3 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) {
+    assert(k % QK_DT3 == 0);
+    const int64_t nb = k / QK_DT3;
+
+    for (int64_t i = 0; i < nb; ++i) {
+        int8_t t1[QK_DT3];
+        int8_t t2[QK_DT3];
+
+        unpack_plane_dt3(x[i].qs[0], x[i].qh[0], t1);
+        unpack_plane_dt3(x[i].qs[1], x[i].qh[1], t2);
+
+        const float d1 = GGML_FP16_TO_FP32(x[i].d[0]);
+        const float d2 = GGML_FP16_TO_FP32(x[i].d[1]);
+
+        for (int j = 0; j < QK_DT3; ++j) {
+            *y++ = d1*t1[j] + d2*t2[j];
+        }
+    }
+}
+
 // ====================== "True" 2-bit (de)-quantization
 
 void dequantize_row_iq2_xxs(const block_iq2_xxs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) {
@@ -5604,6 +5799,38 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte
             {
                 VALIDATE_ROW_DATA_D_F16_IMPL(block_tq2_0, data, nb);
             } break;
+        case GGML_TYPE_DT3:
+            {
+                const block_dt3 * q = (const block_dt3 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    if (!validate_fp16(q[i].d[0], i) || !validate_fp16(q[i].d[1], i)) {
+                        return false;
+                    }
+                    // the ceiling-division packing reaches only 243 of the 256
+                    // byte values in qs, and only 81 in qh (4 trits plus an
+                    // always-zero padding digit); anything else is corruption,
+                    // and a corrupted block would otherwise load and generate
+                    // garbage without any warning
+                    for (int p = 0; p < 2; ++p) {
+                        for (size_t j = 0; j < sizeof(q[i].qs[p]); ++j) {
+                            const uint32_t b = q[i].qs[p][j];
+                            const uint32_t v = (b*243) >> 8; // inverse of the ceiling division
+                            if ((v*256 + 242)/243 != b) {
+                                fprintf(stderr, "ggml_validate_row_data: found invalid dt3 qs byte 0x%02x at block %zu\n", (unsigned) b, i);
+                                return false;
+                            }
+                        }
+                        for (size_t j = 0; j < sizeof(q[i].qh[p]); ++j) {
+                            const uint32_t b = q[i].qh[p][j];
+                            const uint32_t v = (b*243) >> 8;
+                            if (v % 3 != 0 || (v*256 + 242)/243 != b) {
+                                fprintf(stderr, "ggml_validate_row_data: found invalid dt3 qh byte 0x%02x at block %zu\n", (unsigned) b, i);
+                                return false;
+                            }
+                        }
+                    }
+                }
+            } break;
         case GGML_TYPE_IQ1_S:
             {
                 VALIDATE_ROW_DATA_D_F16_IMPL(block_iq1_s, data, nb);

ggml/src/ggml-quants.h

@@ -35,6 +35,7 @@ GGML_API void quantize_row_q8_K_ref(const float * GGML_RESTRICT x, block_q8_K *
 
 GGML_API void quantize_row_tq1_0_ref(const float * GGML_RESTRICT x, block_tq1_0 * GGML_RESTRICT y, int64_t k);
 GGML_API void quantize_row_tq2_0_ref(const float * GGML_RESTRICT x, block_tq2_0 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_dt3_ref  (const float * GGML_RESTRICT x, block_dt3   * GGML_RESTRICT y, int64_t k);
 
 GGML_API void quantize_row_iq3_xxs_ref(const float * GGML_RESTRICT x, block_iq3_xxs * GGML_RESTRICT y, int64_t k);
 GGML_API void quantize_row_iq4_nl_ref (const float * GGML_RESTRICT x, block_iq4_nl  * GGML_RESTRICT y, int64_t k);
@@ -64,6 +65,12 @@ GGML_API void dequantize_row_q8_K(const block_q8_K * GGML_RESTRICT x, float * GG
 
 GGML_API void dequantize_row_tq1_0(const block_tq1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
 GGML_API void dequantize_row_tq2_0(const block_tq2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_dt3  (const block_dt3   * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+
+// unpacks one DT3 plane (128 trits in {-1, 0, 1}) in the original element order;
+// shared between dequantize_row_dt3 and the CPU vec_dot so that the trit
+// decoding exists in exactly one place
+GGML_API void unpack_plane_dt3(const uint8_t * GGML_RESTRICT qs, const uint8_t * GGML_RESTRICT qh, int8_t * GGML_RESTRICT t);
 
 GGML_API void dequantize_row_iq2_xxs(const block_iq2_xxs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
 GGML_API void dequantize_row_iq2_xs (const block_iq2_xs  * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
@@ -88,6 +95,7 @@ GGML_API size_t quantize_iq3_s  (const float * GGML_RESTRICT src, void * GGML_RE
 
 GGML_API size_t quantize_tq1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 GGML_API size_t quantize_tq2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_dt3  (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 
 GGML_API size_t quantize_q2_K(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 GGML_API size_t quantize_q3_K(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);

ggml/src/ggml.c

@@ -924,6 +924,14 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = {
         .to_float                 = (ggml_to_float_t) dequantize_row_tq2_0,
         .from_float_ref           = (ggml_from_float_t) quantize_row_tq2_0_ref,
     },
+    [GGML_TYPE_DT3] = {
+        .type_name                = "dt3",
+        .blck_size                = QK_DT3,
+        .type_size                = sizeof(block_dt3),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_dt3,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_dt3_ref,
+    },
     [36] = { // GGML_TYPE_IQ4_NL_4_4
         .type_name                = "TYPE_IQ4_NL_4_4 REMOVED, use IQ4_NL with runtime repacking",
         .blck_size                = 0,
@@ -1434,6 +1442,7 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) {
         case GGML_FTYPE_MOSTLY_Q4_1:          wtype = GGML_TYPE_Q4_1;  break;
         case GGML_FTYPE_MOSTLY_Q1_0:          wtype = GGML_TYPE_Q1_0;  break;
         case GGML_FTYPE_MOSTLY_Q2_0:          wtype = GGML_TYPE_Q2_0;  break;
+        case GGML_FTYPE_MOSTLY_DT3:           wtype = GGML_TYPE_DT3;   break;
         case GGML_FTYPE_MOSTLY_Q5_0:          wtype = GGML_TYPE_Q5_0;  break;
         case GGML_FTYPE_MOSTLY_Q5_1:          wtype = GGML_TYPE_Q5_1;  break;
         case GGML_FTYPE_MOSTLY_Q8_0:          wtype = GGML_TYPE_Q8_0;  break;
@@ -7957,6 +7966,7 @@ size_t ggml_quantize_chunk(
         case GGML_TYPE_Q6_K:    result = quantize_q6_K   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_TQ1_0:   result = quantize_tq1_0  (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_TQ2_0:   result = quantize_tq2_0  (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_DT3:     result = quantize_dt3    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_IQ2_XXS: result = quantize_iq2_xxs(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_IQ2_XS:  result = quantize_iq2_xs (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_IQ3_XXS: result = quantize_iq3_xxs(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;

gguf-py/gguf/constants.py

@@ -5019,6 +5019,7 @@ class GGMLQuantizationType(IntEnum):
     NVFP4   = 40
     Q1_0    = 41
     Q2_0    = 42
+    DT3     = 43
 
 
 class ExpertGatingFuncType(IntEnum):
@@ -5075,6 +5076,7 @@ class LlamaFileType(IntEnum):
     MOSTLY_NVFP4         = 39  # except 1d tensors
     MOSTLY_Q1_0          = 40  # except 1d tensors
     MOSTLY_Q2_0          = 41  # except 1d tensors
+    MOSTLY_DT3           = 42  # except 1d tensors
 
     GUESSED              = 1024  # not specified in the model file
 
@@ -5206,6 +5208,7 @@ GGML_QUANT_SIZES: dict[GGMLQuantizationType, tuple[int, int]] = {
     GGMLQuantizationType.NVFP4:   (64, 4 + 32),
     GGMLQuantizationType.Q1_0:    (128, 2 + 16),
     GGMLQuantizationType.Q2_0:    (64, 2 + 16),
+    GGMLQuantizationType.DT3:     (128, 2 * (24 + 2 + 2)),
 }
 
 

gguf-py/gguf/quants.py

@@ -654,6 +654,42 @@ class TQ2_0(__Quant, qtype=GGMLQuantizationType.TQ2_0):
         return (d * qs.astype(np.float32))
 
 
+class DT3(__Quant, qtype=GGMLQuantizationType.DT3):
+    # Dual-plane ternary: w = d1 * t1 + d2 * t2 with t in {-1, 0, 1}. Each of
+    # the two planes is packed like TQ1_0 with all constants halved (blocks of
+    # 128 elements): 24 bytes of qs (passes of 16 and 8 bytes) and 2 bytes of
+    # qh per plane, then the two fp16 scales.
+    # Quantization is intentionally not implemented here: DT3 planes are
+    # produced by an external solver (PTQTP) and packed directly.
+    @classmethod
+    def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray:
+        n_blocks = blocks.shape[0]
+
+        qs, rest = np.hsplit(blocks, [2 * 24])
+        qh, d = np.hsplit(rest, [2 * 2])
+
+        d = d.view(np.float16).astype(np.float32).reshape((n_blocks, 2, 1))
+
+        planes = []
+        for p in range(2):
+            pqs = qs[..., p * 24:(p + 1) * 24]
+            pqh = qh[..., p * 2:(p + 1) * 2]
+
+            qs0, qs1 = pqs[..., :16], pqs[..., 16:]
+            qs0 = qs0.reshape((n_blocks, -1, 1, 16)) * np.array([1, 3, 9, 27, 81], dtype=np.uint8).reshape((1, 1, 5, 1))
+            qs0 = qs0.reshape((n_blocks, -1))
+            qs1 = qs1.reshape((n_blocks, -1, 1, 8)) * np.array([1, 3, 9, 27, 81], dtype=np.uint8).reshape((1, 1, 5, 1))
+            qs1 = qs1.reshape((n_blocks, -1))
+            # only 4 trits per qh byte, the 5th base-3 digit is packer padding
+            pqh = pqh.reshape((n_blocks, -1, 1, 2)) * np.array([1, 3, 9, 27], dtype=np.uint8).reshape((1, 1, 4, 1))
+            pqh = pqh.reshape((n_blocks, -1))
+            pq = np.concatenate([qs0, qs1, pqh], axis=-1)
+            pq = ((pq.astype(np.uint16) * 3) >> 8).astype(np.int8) - np.int8(1)
+            planes.append(pq.astype(np.float32))
+
+        return d[:, 0] * planes[0] + d[:, 1] * planes[1]
+
+
 class MXFP4(__Quant, qtype=GGMLQuantizationType.MXFP4):
     # e2m1 values (doubled)
     # ref: https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf

include/llama.h

@@ -156,6 +156,7 @@ extern "C" {
         LLAMA_FTYPE_MOSTLY_NVFP4         = 39, // except 1d tensors
         LLAMA_FTYPE_MOSTLY_Q1_0          = 40, // except 1d tensors
         LLAMA_FTYPE_MOSTLY_Q2_0          = 41, // except 1d tensors
+        LLAMA_FTYPE_MOSTLY_DT3           = 42, // except 1d tensors
 
         LLAMA_FTYPE_GUESSED = 1024, // not specified in the model file
     };

src/llama-model-loader.cpp

@@ -58,6 +58,7 @@ const char * llama_ftype_name(llama_ftype ftype) {
         case LLAMA_FTYPE_MOSTLY_Q6_K:      name = LLAMA_FTYPE_PREFIX "Q6_K"; break;
         case LLAMA_FTYPE_MOSTLY_TQ1_0:     name = LLAMA_FTYPE_PREFIX "TQ1_0 - 1.69 bpw ternary"; break;
         case LLAMA_FTYPE_MOSTLY_TQ2_0:     name = LLAMA_FTYPE_PREFIX "TQ2_0 - 2.06 bpw ternary"; break;
+        case LLAMA_FTYPE_MOSTLY_DT3:       name = LLAMA_FTYPE_PREFIX "DT3 - 3.5 bpw dual-plane ternary"; break;
         case LLAMA_FTYPE_MOSTLY_IQ2_XXS:   name = LLAMA_FTYPE_PREFIX "IQ2_XXS - 2.0625 bpw"; break;
         case LLAMA_FTYPE_MOSTLY_IQ2_XS:    name = LLAMA_FTYPE_PREFIX "IQ2_XS - 2.3125 bpw"; break;
         case LLAMA_FTYPE_MOSTLY_IQ2_S:     name = LLAMA_FTYPE_PREFIX "IQ2_S - 2.5 bpw"; break;
@@ -747,6 +748,7 @@ llama_model_loader::llama_model_loader(
             case GGML_TYPE_Q6_K:    ftype = LLAMA_FTYPE_MOSTLY_Q6_K;    break;
             case GGML_TYPE_TQ1_0:   ftype = LLAMA_FTYPE_MOSTLY_TQ1_0;   break;
             case GGML_TYPE_TQ2_0:   ftype = LLAMA_FTYPE_MOSTLY_TQ2_0;   break;
+            case GGML_TYPE_DT3:     ftype = LLAMA_FTYPE_MOSTLY_DT3;     break;
             case GGML_TYPE_IQ2_XXS: ftype = LLAMA_FTYPE_MOSTLY_IQ2_XXS; break;
             case GGML_TYPE_IQ2_XS:  ftype = LLAMA_FTYPE_MOSTLY_IQ2_XS;  break;
             case GGML_TYPE_IQ2_S:   ftype = LLAMA_FTYPE_MOSTLY_IQ2_S;   break;

src/llama-quant.cpp

@@ -396,7 +396,8 @@ static ggml_type tensor_type_fallback(quantize_state_impl & qs, const ggml_tenso
             case GGML_TYPE_Q2_K:
             case GGML_TYPE_Q3_K:
             case GGML_TYPE_TQ1_0:
-            case GGML_TYPE_TQ2_0:   return_type = GGML_TYPE_Q4_0;   break;
+            case GGML_TYPE_TQ2_0:
+            case GGML_TYPE_DT3:     return_type = GGML_TYPE_Q4_0;   break;
             case GGML_TYPE_Q4_K:    return_type = GGML_TYPE_Q5_0;   break;
             case GGML_TYPE_Q5_K:    return_type = GGML_TYPE_Q5_1;   break;
             case GGML_TYPE_Q6_K:    return_type = GGML_TYPE_Q8_0;   break;
@@ -493,7 +494,7 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type
             else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS) {
                 new_type = GGML_TYPE_IQ3_S;
             }
-            else if (ftype == LLAMA_FTYPE_MOSTLY_TQ1_0 || ftype == LLAMA_FTYPE_MOSTLY_TQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_Q2_0) {
+            else if (ftype == LLAMA_FTYPE_MOSTLY_TQ1_0 || ftype == LLAMA_FTYPE_MOSTLY_TQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_Q2_0 || ftype == LLAMA_FTYPE_MOSTLY_DT3) {
                 new_type = GGML_TYPE_Q4_K;
             }
         }
@@ -831,6 +832,7 @@ ggml_type llama_ftype_get_default_type(llama_ftype ftype) {
         case LLAMA_FTYPE_MOSTLY_Q6_K:    return GGML_TYPE_Q6_K;
         case LLAMA_FTYPE_MOSTLY_TQ1_0:   return GGML_TYPE_TQ1_0;
         case LLAMA_FTYPE_MOSTLY_TQ2_0:   return GGML_TYPE_TQ2_0;
+        case LLAMA_FTYPE_MOSTLY_DT3:     return GGML_TYPE_DT3;
         case LLAMA_FTYPE_MOSTLY_IQ2_XXS: return GGML_TYPE_IQ2_XXS;
         case LLAMA_FTYPE_MOSTLY_IQ2_XS:  return GGML_TYPE_IQ2_XS;
         case LLAMA_FTYPE_MOSTLY_IQ2_S:   return GGML_TYPE_IQ2_XS;
@@ -882,6 +884,11 @@ static void llama_model_quantize_impl(const std::string & fname_inp, const std::
         throw std::runtime_error(format("invalid output file type %d\n", ftype));
     }
 
+    if (ftype == LLAMA_FTYPE_MOSTLY_DT3) {
+        LLAMA_LOG_WARN("%s: DT3 is being quantized with the in-tree REFERENCE quantizer only - "
+                       "the measured DT3 quality requires the planes from ternaria's PTQTP pipeline\n", __func__);
+    }
+
     // mmap consistently increases speed on Linux, and also increases speed on Windows with
     // hot cache. It may cause a slowdown on macOS, possibly related to free memory.
 #if defined(__linux__) || defined(_WIN32)

tests/CMakeLists.txt

@@ -289,6 +289,8 @@ if (NOT GGML_BACKEND_DL)
     # these tests use the backends directly and cannot be built with dynamic loading
     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,527 @@
+// 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.
+//
+// Errors are judged on the relative Frobenius norm, ||gpu - ref|| / ||ref||;
+// the elementwise maximum is reported as information only, since it explodes
+// on cancellation whenever a true output value is near zero.
+//
+// MUL_MAT with more destination columns than the MMVQ limit falls back to
+// dequantization + cuBLAS GEMM, which on fast-fp16 hardware rounds the
+// dequantized weights to fp16. DT3 weights (d1*t1 + d2*t2, the sum of two
+// fp16-scaled terms) are generally NOT fp16-representable, so that path is
+// judged against a reference computed from fp16-rounded weights (taking the
+// better of the two references, so the test also passes when
+// GGML_CUDA_CUBLAS_COMPUTE_TYPE=f32 disables the rounding). Q4_1 (same
+// regime: d*q + m not fp16-exact) and Q4_0 (weights fp16-exact) go through
+// the identical comparison as controls, reported but gated loosely.
+//
+// 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.
+//
+// DT3 is implemented for CUDA and HIP only. Without one of those backends the
+// test is skipped and succeeds — an unsupported backend is not a failure.
+
+#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;
+}
+
+struct mat_err {
+    double norm_rel; // ||gpu - ref|| / ||ref||
+    double max_rel;  // max elementwise |gpu - ref| / max(|ref|, 1) — information only
+};
+
+static mat_err compare_mat(const std::vector<float> & gpu, const std::vector<double> & ref) {
+    double num  = 0.0;
+    double den  = 0.0;
+    double mrel = 0.0;
+    for (size_t i = 0; i < gpu.size(); ++i) {
+        const double diff = (double)gpu[i] - ref[i];
+        num += diff*diff;
+        den += ref[i]*ref[i];
+        const double rel = fabs(diff) / (fabs(ref[i]) > 1.0 ? fabs(ref[i]) : 1.0);
+        mrel = rel > mrel ? rel : mrel;
+    }
+    return { sqrt(num/den), mrel };
+}
+
+// MUL_MAT on the GPU vs double precision references from the dequantized
+// weights (exact, and rounded to fp16 as the GEMM fallback does).
+// strict = tight gates (DT3); controls are gated loosely at 1e-2.
+static int test_mul_mat(ggml_backend_t backend, ggml_type type, const std::vector<uint8_t> & data,
+                        const std::vector<float> & ref_w, const std::vector<float> & y, bool strict) {
+    int num_failed = 0;
+
+    const int ncols_dst[] = {1, 2, 5, 8, 16};
+
+    // the same weights as the fp16 GEMM fallback sees them
+    std::vector<float> ref_w16(ref_w.size());
+    for (size_t i = 0; i < ref_w.size(); ++i) {
+        ref_w16[i] = ggml_fp16_to_fp32(ggml_fp32_to_fp16(ref_w[i]));
+    }
+
+    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, type, 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 %s\n", ggml_type_name(type));
+            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);
+
+        // references in double from the exact and the fp16-rounded weights
+        std::vector<double> ref((size_t)NROWS*n);
+        std::vector<double> ref16((size_t)NROWS*n);
+        for (int j = 0; j < n; ++j) {
+            for (int r = 0; r < NROWS; ++r) {
+                double sum   = 0.0;
+                double sum16 = 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];
+                    sum16 += (double)ref_w16[(size_t)r*NCOLS + k] * (double)y[(size_t)j*NCOLS + k];
+                }
+                ref  [(size_t)j*NROWS + r] = sum;
+                ref16[(size_t)j*NROWS + r] = sum16;
+            }
+        }
+
+        const mat_err err   = compare_mat(gpu, ref);
+        const mat_err err16 = compare_mat(gpu, ref16);
+
+        // n <= 8 is the MMVQ path with exact integer dot products, judged
+        // against the exact reference. Larger n is the dequantize + GEMM
+        // fallback whose numerics (fp16 or TF32 compute, depending on the
+        // hardware and on GGML_CUDA_CUBLAS_COMPUTE_TYPE) are cuBLAS's, not
+        // ours: for the strict type it is gated below by bit-identity with
+        // the same GEMM on an F16 tensor, and only reported here.
+        // This mirrors MMVQ_MAX_BATCH_SIZE (8) from ggml-cuda/mmvq.cu by hand,
+        // because the constant and the per-arch should_use_mmvq tables are not
+        // exported. If upstream raises the limit, or an architecture routes a
+        // larger batch through MMVQ, this gating goes stale silently: n = 16
+        // would take the MMVQ path but still be judged as the GEMM one, which
+        // only loosens the check, never tightens it. Whoever touches the MMVQ
+        // dispatch should revisit this line.
+        const bool   is_mmvq = n <= 8;
+        const bool   gated   = is_mmvq || !strict;
+        const double err_gate = is_mmvq ? err.norm_rel : (err.norm_rel < err16.norm_rel ? err.norm_rel : err16.norm_rel);
+        const double tol      = strict ? 1e-5 : 1e-2;
+        const bool   failed   = gated && err_gate > tol;
+        printf("%s: %s mul_mat GPU, ncols_dst = %2d (%s): norm rel err vs exact ref = %g, vs fp16 ref = %g (max elem rel: %g)\n",
+               failed ? "FAILED" : gated ? "OK" : "INFO", ggml_type_name(type), n, is_mmvq ? "MMVQ" : "GEMM",
+               err.norm_rel, err16.norm_rel, err.max_rel);
+        if (failed) {
+            num_failed++;
+        }
+
+        ggml_backend_buffer_free(buf);
+        ggml_free(ctx);
+    }
+
+    // MMVQ vs the dequantization-based path: the first 8 columns of the GEMM
+    // run must match the ncols_dst = 8 MMVQ run to fp16 weight rounding
+    {
+        const std::vector<float> & mmvq = results[3]; // n = 8
+        const std::vector<float> & gemm = results[4]; // n = 16
+        double num = 0.0;
+        double den = 0.0;
+        for (int j = 0; j < 8; ++j) {
+            for (int r = 0; r < NROWS; ++r) {
+                const double diff = (double)mmvq[(size_t)j*NROWS + r] - (double)gemm[(size_t)j*NROWS + r];
+                num += diff*diff;
+                den += (double)mmvq[(size_t)j*NROWS + r]*(double)mmvq[(size_t)j*NROWS + r];
+            }
+        }
+        const double norm_rel = sqrt(num/den);
+        const double tol      = strict ? 5e-3 : 1e-2;
+        printf("%s: %s MMVQ vs GEMM path on shared columns: norm rel err = %g\n",
+               norm_rel <= tol ? "OK" : "FAILED", ggml_type_name(type), norm_rel);
+        if (norm_rel > tol) {
+            num_failed++;
+        }
+    }
+
+    // the GEMM fallback must be exactly "as if the weights were an F16
+    // tensor holding fp16(dequant(block))": running the same GEMM with an
+    // F16 src0 built from the fp16-rounded reference weights must give a
+    // bit-identical result. This isolates our (already bit-validated)
+    // dequantization from cuBLAS numerics.
+    if (strict) {
+        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 * a16 = ggml_new_tensor_2d(ctx, GGML_TYPE_F16, NCOLS, NROWS);
+        ggml_tensor * b   = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, NCOLS, 16);
+        ggml_tensor * out = ggml_mul_mat(ctx, a16, b);
+
+        ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, backend);
+        GGML_ASSERT(buf != nullptr);
+
+        std::vector<ggml_fp16_t> w16(ref_w.size());
+        for (size_t i = 0; i < ref_w.size(); ++i) {
+            w16[i] = ggml_fp32_to_fp16(ref_w[i]);
+        }
+        ggml_backend_tensor_set(a16, w16.data(), 0, w16.size()*sizeof(ggml_fp16_t));
+        ggml_backend_tensor_set(b,   y.data(),   0, (size_t)NCOLS*16*sizeof(float));
+
+        std::vector<float> gpu16((size_t)NROWS*16);
+        compute_graph(backend, ctx, out, gpu16.data());
+
+        const std::vector<float> & gemm = results[4]; // n = 16
+        int    n_mismatch = 0;
+        double max_diff   = 0.0;
+        for (size_t i = 0; i < gemm.size(); ++i) {
+            const double diff = fabs((double)gemm[i] - (double)gpu16[i]);
+            max_diff = diff > max_diff ? diff : max_diff;
+            if (gemm[i] != gpu16[i]) {
+                n_mismatch++;
+            }
+        }
+        printf("%s: %s GEMM path vs F16 GEMM on fp16-rounded weights: %d mismatches, max |diff| = %g\n",
+               n_mismatch == 0 ? "OK" : "FAILED", ggml_type_name(type), n_mismatch, max_diff);
+        if (n_mismatch != 0) {
+            num_failed++;
+        }
+
+        ggml_backend_buffer_free(buf);
+        ggml_free(ctx);
+    }
+
+    // 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
+    if (type == GGML_TYPE_DT3) {
+        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;
+}
+
+// quantize random floats to a control type and return raw data + dequantized
+// reference weights
+static void build_control_data(ggml_type type, std::vector<uint8_t> & data, std::vector<float> & ref_w) {
+    std::vector<float> src((size_t)NROWS*NCOLS);
+    for (size_t i = 0; i < src.size(); ++i) {
+        src[i] = ((int)(rng_next() % 2001) - 1000)/1000.0f;
+    }
+
+    data.resize(ggml_row_size(type, NCOLS)*NROWS);
+    const size_t written = ggml_quantize_chunk(type, src.data(), data.data(), 0, NROWS, NCOLS, nullptr);
+    GGML_ASSERT(written == data.size());
+
+    ref_w.resize(src.size());
+    ggml_get_type_traits(type)->to_float(data.data(), ref_w.data(), (int64_t)NROWS*NCOLS);
+}
+
+int main(void) {
+    // Only CUDA and HIP (which reports itself as "ROCm") implement DT3. Any
+    // other GPU backend is skipped rather than failed: Vulkan and SYCL answer
+    // supports_op == false for DT3, which is the correct answer for them and
+    // not a bug to report, and Metal answers true for almost any type but has
+    // no DT3 shader, so it would die in pipeline compilation mid-test. Picking
+    // the backend by name keeps this test honest on machines we do not have.
+    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) {
+            continue;
+        }
+        const char * name = ggml_backend_dev_name(dev);
+        if (strncmp(name, "CUDA", 4) != 0 && strncmp(name, "ROCm", 4) != 0) {
+            printf("skipping GPU backend %s: DT3 is only implemented for CUDA/HIP\n", name);
+            continue;
+        }
+        backend = ggml_backend_dev_init(dev, nullptr);
+        printf("using GPU backend: %s\n", name);
+        break;
+    }
+    if (backend == nullptr) {
+        printf("no CUDA/HIP 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);
+
+    // activations: integers with amax 127 in every 32-element chunk of every
+    // column, so their q8_1 quantization is exact
+    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);
+    }
+
+    int num_failed = 0;
+    num_failed += test_dequant(backend, data, ref);
+    num_failed += test_mul_mat(backend, GGML_TYPE_DT3, data, ref, y, /*strict =*/ true);
+
+    // controls through the identical comparison: Q4_1 shares DT3's regime
+    // (dequantized weights not fp16-exact), Q4_0's weights are fp16-exact
+    // and show the pure GEMM error floor
+    for (ggml_type control : {GGML_TYPE_Q4_1, GGML_TYPE_Q4_0}) {
+        std::vector<uint8_t> cdata;
+        std::vector<float>   cref;
+        build_control_data(control, cdata, cref);
+        num_failed += test_mul_mat(backend, control, cdata, cref, y, /*strict =*/ false);
+    }
+
+    ggml_backend_free(backend);
+
+    if (num_failed > 0) {
+        printf("%d tests FAILED\n", num_failed);
+        return 1;
+    }
+    printf("all tests OK\n");
+    return 0;
+}

tests/test-dt3-rust-parity.py

@@ -0,0 +1,79 @@
+#!/usr/bin/env python3
+# Parity check between ternaria's Rust DT3 packer and llama.cpp's C
+# dequantization: blocks packed by pack_dt3 must dequantize (through
+# dequantize_row_dt3, exposed by `test-dt3 --dequant`) to exactly
+# d1*t1 + d2*t2 computed independently in numpy from the same trits and
+# fp16-rounded scales.
+#
+# Must run inside the ternaria environment, e.g.:
+#   cd /path/to/ternaria && uv run python /path/to/llama.cpp/tests/test-dt3-rust-parity.py \
+#       /path/to/llama.cpp/build/bin/test-dt3
+
+import subprocess
+import sys
+import tempfile
+from pathlib import Path
+
+import numpy as np
+
+from ternaria._core import pack_dt3
+
+QK_DT3 = 128
+BLOCK_BYTES = 56
+
+
+def main() -> int:
+    if len(sys.argv) != 2:
+        print(f"usage: {sys.argv[0]} /path/to/test-dt3", file=sys.stderr)
+        return 1
+    test_bin = Path(sys.argv[1])
+
+    rng = np.random.default_rng(20260810)
+
+    rows, cols = 8, 1024
+    t1 = rng.integers(-1, 2, size=(rows, cols)).astype(np.int8)
+    t2 = rng.integers(-1, 2, size=(rows, cols)).astype(np.int8)
+    s1 = rng.normal(size=(rows, cols // QK_DT3)).astype(np.float32)
+    s2 = (0.25 * rng.normal(size=(rows, cols // QK_DT3))).astype(np.float32)
+
+    # edge cases: an all-zero block, a block with negative scales, and a
+    # block that is non-trivial only in the qh region (elements 120..127)
+    t1[0, :QK_DT3] = 0
+    t2[0, :QK_DT3] = 0
+    s1[0, 0] = 0.0
+    s2[0, 0] = 0.0
+    s1[0, 1] = -abs(s1[0, 1])
+    s2[0, 1] = -abs(s2[0, 1])
+    t1[1, :QK_DT3] = 0
+    t2[1, :QK_DT3] = 0
+    t1[1, 120:128] = [-1, 1, 0, -1, 1, -1, 0, 1]
+    t2[1, 120:128] = [1, -1, 1, 0, 0, 1, -1, -1]
+
+    raw = np.asarray(pack_dt3(t1, s1, t2, s2), dtype=np.uint8)
+    assert raw.size == rows * (cols // QK_DT3) * BLOCK_BYTES, raw.size
+
+    with tempfile.TemporaryDirectory() as tmp:
+        raw_path = Path(tmp) / "dt3.bin"
+        out_path = Path(tmp) / "out.f32"
+        raw_path.write_bytes(raw.tobytes())
+        subprocess.run([str(test_bin), "--dequant", str(raw_path), str(out_path)], check=True)
+        got = np.fromfile(out_path, dtype=np.float32).reshape(rows, cols)
+
+    # what the packed bytes mean: fp16-rounded scales times the trits
+    d1 = s1.astype(np.float16).astype(np.float32).repeat(QK_DT3, axis=1)
+    d2 = s2.astype(np.float16).astype(np.float32).repeat(QK_DT3, axis=1)
+    expected = d1 * t1.astype(np.float32) + d2 * t2.astype(np.float32)
+
+    if not np.array_equal(got, expected):
+        bad = np.nonzero(got != expected)
+        print(f"FAILED: {len(bad[0])} of {got.size} elements differ", file=sys.stderr)
+        r, c = bad[0][0], bad[1][0]
+        print(f"first mismatch at ({r}, {c}): got {got[r, c]}, expected {expected[r, c]}", file=sys.stderr)
+        return 1
+
+    print(f"ok: {got.size} weights bit-exact between Rust pack_dt3 and C dequantize_row_dt3")
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())

tests/test-dt3.cpp

@@ -0,0 +1,648 @@
+// Bit-level unit tests for the DT3 dual-plane ternary format
+//
+// DT3 packs 128 weights as two ternary planes (w = d1*t1 + d2*t2), each plane
+// laid out exactly like tq1_0 with all constants halved. A block with
+// misplaced bits still loads and generates, so the layout is checked here
+// bit by bit against an independent packer that implements the format
+// specification directly.
+//
+// Extra mode for cross-implementation parity checks (see
+// tests/test-dt3-rust-parity.py):
+//   test-dt3 --dequant IN.bin OUT.f32
+// dequantizes raw DT3 blocks from IN.bin into float32 little-endian OUT.f32.
+
+#include "ggml.h"
+#include "ggml-cpu.h"
+
+#undef NDEBUG
+#include <assert.h>
+#include <inttypes.h>
+#include <math.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.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 (not from
+// ggml-quants.c): trits in {-1, 0, 1}, element i of the 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)
+// with the first element in the most significant trit, an extra *3 shift in
+// region C, and ceiling division by 243 to fit 5 trits per byte.
+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 = 0x12345678;
+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;
+}
+
+static int test_layout_constants(void) {
+    int num_failed = 0;
+
+    if (ggml_blck_size(GGML_TYPE_DT3) != QK_DT3) {
+        printf("FAILED: blck_size is %" PRId64 ", expected %d\n", ggml_blck_size(GGML_TYPE_DT3), QK_DT3);
+        num_failed++;
+    }
+    if (ggml_type_size(GGML_TYPE_DT3) != DT3_BLOCK_SIZE) {
+        printf("FAILED: type_size is %zu, expected %zu\n", ggml_type_size(GGML_TYPE_DT3), DT3_BLOCK_SIZE);
+        num_failed++;
+    }
+    // 56 bytes / 128 weights = 3.5 bpw exactly
+    if (DT3_BLOCK_SIZE*8 != (size_t)QK_DT3*7/2) {
+        printf("FAILED: not 3.5 bpw\n");
+        num_failed++;
+    }
+
+    return num_failed;
+}
+
+// a single trit set to -1 or +1 at each position of each plane must come back
+// at the same position, scaled by that plane's scale only
+static int test_single_trits(const ggml_type_traits * qfns) {
+    int num_failed = 0;
+
+    const float d1 = 1.0f;   // exact in fp16
+    const float d2 = 0.25f;  // exact in fp16
+
+    for (int plane = 0; plane < 2; ++plane) {
+        for (int pos = 0; pos < QK_DT3; ++pos) {
+            for (int val = -1; val <= 1; val += 2) {
+                int8_t t1[QK_DT3] = {0};
+                int8_t t2[QK_DT3] = {0};
+                (plane == 0 ? t1 : t2)[pos] = (int8_t)val;
+
+                uint8_t block[DT3_BLOCK_SIZE];
+                ref_pack_block(t1, d1, t2, d2, block);
+
+                float out[QK_DT3];
+                qfns->to_float(block, out, QK_DT3);
+
+                for (int j = 0; j < QK_DT3; ++j) {
+                    const float expected = j == pos ? (plane == 0 ? d1 : d2)*val : 0.0f;
+                    if (out[j] != expected) {
+                        printf("FAILED: plane %d pos %d val %d: out[%d] = %f, expected %f\n",
+                               plane, pos, val, j, out[j], expected);
+                        num_failed++;
+                    }
+                }
+            }
+        }
+    }
+
+    return num_failed;
+}
+
+// an all-zero plane packs to qs = 128, qh = 127 everywhere (hand-computed:
+// qs: xi = 1 for the 5 trits -> q = 121 -> ceil(121*256/243) = 128,
+// qh: q = 40*3 = 120 -> ceil(120*256/243) = 127)
+constexpr uint8_t ZERO_QS = 128;
+constexpr uint8_t ZERO_QH = 127;
+
+static void build_zero_pattern_block(uint8_t * block, float d1, float d2) {
+    memset(block + OFF_QS, ZERO_QS, 2*DT3_QS_BYTES);
+    memset(block + OFF_QH, ZERO_QH, 2*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));
+}
+
+// check the byte positions of the region boundaries against literal values
+// computed by hand from the packing formula (and cross-checked one by one
+// against an independent implementation of the spec). Both directions go
+// against the LIBRARY code — to_float and from_float — never against a
+// packer defined in this file, so a bit-placement bug in ggml-quants.c
+// cannot cancel out.
+static int test_byte_positions(const ggml_type_traits * qfns, const ggml_type_traits_cpu * qfns_cpu) {
+    int num_failed = 0;
+
+    struct byte_case {
+        int     pos;   // element with t = -1, all others 0
+        size_t  off;   // byte offset inside the plane data (qs: 0..23, qh: 24, 25)
+        uint8_t value; // hand-computed byte value at that offset
+    };
+    const byte_case cases[] = {
+        {   0,  0,  43 }, // region A: most significant trit of qs[0]:  q = 121-81 = 40  -> 43
+        {  16,  0, 100 }, // region A: second trit of qs[0]:            q = 121-27 = 94  -> 100
+        {  79, 15, 127 }, // region A boundary: last trit of qs[15]:    q = 121-1  = 120 -> 127
+        {  80, 16,  43 }, // region B boundary: first trit of qs[16]:   q = 40           -> 43
+        { 119, 23, 127 }, // region B boundary: last trit of qs[23]:    q = 120          -> 127
+        { 120, 24,  42 }, // region C boundary: first trit of qh[0]:    q = (40-27)*3    -> 42
+        { 127, 25, 124 }, // region C: last element, qh[1]:             q = (40-1)*3     -> 124
+    };
+
+    for (size_t c = 0; c < sizeof(cases)/sizeof(cases[0]); ++c) {
+        const int     pos = cases[c].pos;
+        const size_t  off = cases[c].off;
+        const uint8_t val = cases[c].value;
+
+        // unpack direction: writing the literal byte must put -1 at exactly
+        // element `pos` of the corresponding plane
+        for (int plane = 0; plane < 2; ++plane) {
+            uint8_t block[DT3_BLOCK_SIZE];
+            build_zero_pattern_block(block, 1.0f, 0.25f);
+            if (off < DT3_QS_BYTES) {
+                block[OFF_QS + plane*DT3_QS_BYTES + off] = val;
+            } else {
+                block[OFF_QH + plane*DT3_QH_BYTES + (off - DT3_QS_BYTES)] = val;
+            }
+
+            float out[QK_DT3];
+            qfns->to_float(block, out, QK_DT3);
+
+            const float dp = plane == 0 ? 1.0f : 0.25f;
+            for (int j = 0; j < QK_DT3; ++j) {
+                const float expected = j == pos ? -dp : 0.0f;
+                if (out[j] != expected) {
+                    printf("FAILED: byte 0x%02x at plane %d offset %zu: out[%d] = %f, expected %f\n",
+                           val, plane, off, j, out[j], expected);
+                    num_failed++;
+                }
+            }
+        }
+
+        // pack direction: quantizing a single -1 must produce exactly the
+        // literal byte at the right offset of plane 0 (plane 1 is the zero
+        // pattern since there is no residual), and nothing else
+        {
+            float x[QK_DT3] = {0.0f};
+            x[pos] = -1.0f;
+
+            uint8_t block[DT3_BLOCK_SIZE];
+            qfns_cpu->from_float(x, block, QK_DT3);
+
+            uint8_t expected[DT3_BLOCK_SIZE];
+            build_zero_pattern_block(expected, 1.0f, 0.0f);
+            if (off < DT3_QS_BYTES) {
+                expected[OFF_QS + off] = val;
+            } else {
+                expected[OFF_QH + (off - DT3_QS_BYTES)] = val;
+            }
+
+            for (size_t b = 0; b < DT3_BLOCK_SIZE; ++b) {
+                if (block[b] != expected[b]) {
+                    printf("FAILED: pack of -1 at element %d: byte %zu is 0x%02x, expected 0x%02x\n",
+                           pos, b, block[b], expected[b]);
+                    num_failed++;
+                }
+            }
+        }
+    }
+
+    // all-zero input packs to the all-zero pattern with zero scales
+    {
+        const float x[QK_DT3] = {0.0f};
+
+        uint8_t block[DT3_BLOCK_SIZE];
+        qfns_cpu->from_float(x, block, QK_DT3);
+
+        uint8_t expected[DT3_BLOCK_SIZE];
+        build_zero_pattern_block(expected, 0.0f, 0.0f);
+
+        for (size_t b = 0; b < DT3_BLOCK_SIZE; ++b) {
+            if (block[b] != expected[b]) {
+                printf("FAILED: pack of all-zero block: byte %zu is 0x%02x, expected 0x%02x\n",
+                       b, block[b], expected[b]);
+                num_failed++;
+            }
+        }
+    }
+
+    return num_failed;
+}
+
+// ggml_validate_row_data must accept exactly the reachable byte values:
+// 243 of 256 in qs, 81 of 256 in qh (4 trits + an always-zero padding digit)
+static int test_validate(void) {
+    int num_failed = 0;
+
+    bool valid_qs[256] = {false};
+    bool valid_qh[256] = {false};
+    for (uint32_t q = 0; q < 243; ++q) {
+        valid_qs[(q*256 + 242)/243] = true;
+        if (q % 3 == 0 && q <= 240) {
+            valid_qh[(q*256 + 242)/243] = true;
+        }
+    }
+
+    printf("(invalid-byte messages below are expected, the validator is being probed)\n");
+
+    for (int b = 0; b < 256; ++b) {
+        uint8_t block[DT3_BLOCK_SIZE];
+
+        build_zero_pattern_block(block, 1.0f, 0.25f);
+        block[OFF_QS] = (uint8_t) b;
+        if (ggml_validate_row_data(GGML_TYPE_DT3, block, DT3_BLOCK_SIZE) != valid_qs[b]) {
+            printf("FAILED: validate qs byte 0x%02x: expected %s\n", b, valid_qs[b] ? "valid" : "invalid");
+            num_failed++;
+        }
+
+        build_zero_pattern_block(block, 1.0f, 0.25f);
+        block[OFF_QH + DT3_QH_BYTES] = (uint8_t) b; // plane 1 qh, to also cover the second plane
+        if (ggml_validate_row_data(GGML_TYPE_DT3, block, DT3_BLOCK_SIZE) != valid_qh[b]) {
+            printf("FAILED: validate qh byte 0x%02x: expected %s\n", b, valid_qh[b] ? "valid" : "invalid");
+            num_failed++;
+        }
+    }
+
+    // a block filled with 0xaa must be rejected (0xaa is a reachable qs code
+    // but violates the qh padding-digit invariant)
+    {
+        uint8_t block[DT3_BLOCK_SIZE];
+        memset(block, 0xaa, DT3_BLOCK_SIZE);
+        const uint16_t h = ggml_fp32_to_fp16(1.0f);
+        memcpy(block + OFF_D,     &h, sizeof(h));
+        memcpy(block + OFF_D + 2, &h, sizeof(h));
+        if (ggml_validate_row_data(GGML_TYPE_DT3, block, DT3_BLOCK_SIZE)) {
+            printf("FAILED: validate accepted a block filled with 0xaa\n");
+            num_failed++;
+        }
+    }
+
+    // packed real data must pass
+    {
+        int8_t t1[QK_DT3];
+        int8_t t2[QK_DT3];
+        for (int j = 0; j < QK_DT3; ++j) {
+            t1[j] = rng_trit();
+            t2[j] = rng_trit();
+        }
+        uint8_t block[DT3_BLOCK_SIZE];
+        ref_pack_block(t1, -0.5f, t2, 0.125f, block);
+        if (!ggml_validate_row_data(GGML_TYPE_DT3, block, DT3_BLOCK_SIZE)) {
+            printf("FAILED: validate rejected a well-formed block\n");
+            num_failed++;
+        }
+    }
+
+    return num_failed;
+}
+
+// random trits and scales (negative scales included) must round-trip exactly
+static int test_roundtrip(const ggml_type_traits * qfns) {
+    int num_failed = 0;
+
+    const float scales[][2] = {
+        { 1.0f,  0.25f  },
+        { 0.5f, -0.125f }, // negative second plane
+        {-2.0f,  0.75f  }, // negative first plane
+        { 0.0f,  0.0f   }, // all-zero scales
+    };
+
+    for (size_t sc = 0; sc < sizeof(scales)/sizeof(scales[0]); ++sc) {
+        for (int rep = 0; rep < 64; ++rep) {
+            int8_t t1[QK_DT3];
+            int8_t t2[QK_DT3];
+            for (int j = 0; j < QK_DT3; ++j) {
+                t1[j] = rng_trit();
+                t2[j] = rng_trit();
+            }
+
+            const float d1 = ggml_fp16_to_fp32(ggml_fp32_to_fp16(scales[sc][0]));
+            const float d2 = ggml_fp16_to_fp32(ggml_fp32_to_fp16(scales[sc][1]));
+
+            uint8_t block[DT3_BLOCK_SIZE];
+            ref_pack_block(t1, d1, t2, d2, block);
+
+            float out[QK_DT3];
+            qfns->to_float(block, out, QK_DT3);
+
+            for (int j = 0; j < QK_DT3; ++j) {
+                const float expected = d1*t1[j] + d2*t2[j];
+                if (out[j] != expected) {
+                    printf("FAILED: roundtrip scales (%f, %f) rep %d: out[%d] = %f, expected %f\n",
+                           d1, d2, rep, j, out[j], expected);
+                    num_failed++;
+                }
+            }
+        }
+    }
+
+    return num_failed;
+}
+
+// the in-tree quantizer must produce the same bytes as the independent packer
+// when the input is already exactly ternary (plane 1 = input, plane 2 = 0)
+static int test_quantize_pack_parity(const ggml_type_traits_cpu * qfns_cpu) {
+    int num_failed = 0;
+
+    for (int rep = 0; rep < 64; ++rep) {
+        int8_t t1[QK_DT3];
+        const int8_t t2[QK_DT3] = {0};
+        float  x[QK_DT3];
+        for (int j = 0; j < QK_DT3; ++j) {
+            t1[j] = rng_trit();
+            x[j]  = (float) t1[j];
+        }
+        // make sure the block is not all zeros so that d1 == 1.0
+        t1[0] = 1;
+        x[0]  = 1.0f;
+
+        uint8_t expected[DT3_BLOCK_SIZE];
+        ref_pack_block(t1, 1.0f, t2, 0.0f, expected);
+
+        uint8_t block[DT3_BLOCK_SIZE];
+        qfns_cpu->from_float(x, block, QK_DT3);
+
+        if (memcmp(block, expected, DT3_BLOCK_SIZE) != 0) {
+            for (size_t b = 0; b < DT3_BLOCK_SIZE; ++b) {
+                if (block[b] != expected[b]) {
+                    printf("FAILED: quantize pack parity rep %d: byte %zu is 0x%02x, expected 0x%02x\n",
+                           rep, b, block[b], expected[b]);
+                }
+            }
+            num_failed++;
+        }
+    }
+
+    return num_failed;
+}
+
+// vec_dot against a hand-made sum over the KNOWN trits (not against our own
+// dequantization): sum_i y_i * (d1*t1_i + d2*t2_i) with y from q8_0's own
+// to_float. Random trits make the qh bytes non-trivial, which would expose a
+// vectorization that reads the padding 5th trit of the qh bytes.
+static int test_vec_dot(const ggml_type_traits_cpu * qfns_cpu) {
+    int num_failed = 0;
+
+    const auto * vdot_traits = ggml_get_type_traits_cpu(qfns_cpu->vec_dot_type);
+    const auto * vdot_qfns   = ggml_get_type_traits(qfns_cpu->vec_dot_type);
+
+    if (qfns_cpu->vec_dot_type != GGML_TYPE_Q8_0) {
+        printf("FAILED: vec_dot_type is %s, expected q8_0\n", ggml_type_name(qfns_cpu->vec_dot_type));
+        return 1;
+    }
+
+    const int nblocks = 4;
+    const int n       = nblocks*QK_DT3;
+
+    for (int rep = 0; rep < 64; ++rep) {
+        std::vector<int8_t> t1(n);
+        std::vector<int8_t> t2(n);
+        std::vector<float>  d1(nblocks);
+        std::vector<float>  d2(nblocks);
+        std::vector<uint8_t> xq(nblocks*DT3_BLOCK_SIZE);
+
+        for (int i = 0; i < nblocks; ++i) {
+            for (int j = 0; j < QK_DT3; ++j) {
+                t1[i*QK_DT3 + j] = rng_trit();
+                t2[i*QK_DT3 + j] = rng_trit();
+            }
+            // fp16-exact scales of both signs
+            d1[i] = (float)((int)(rng_next() % 9) - 4) * 0.25f;
+            d2[i] = (float)((int)(rng_next() % 9) - 4) * 0.0625f;
+            ref_pack_block(&t1[i*QK_DT3], d1[i], &t2[i*QK_DT3], d2[i], &xq[i*DT3_BLOCK_SIZE]);
+        }
+
+        std::vector<float> y(n);
+        for (int j = 0; j < n; ++j) {
+            y[j] = 0.1f + 2.0f*cosf((float)(j + rep));
+        }
+
+        std::vector<uint8_t> yq(ggml_row_size(qfns_cpu->vec_dot_type, n));
+        vdot_traits->from_float(y.data(), yq.data(), n);
+
+        // exact values the integer path sees
+        std::vector<float> ydq(n);
+        vdot_qfns->to_float(yq.data(), ydq.data(), n);
+
+        double ref = 0.0;
+        for (int i = 0; i < nblocks; ++i) {
+            for (int j = 0; j < QK_DT3; ++j) {
+                const int ij = i*QK_DT3 + j;
+                ref += (double)ydq[ij] * ((double)d1[i]*t1[ij] + (double)d2[i]*t2[ij]);
+            }
+        }
+
+        float result = INFINITY;
+        qfns_cpu->vec_dot(n, &result, 0, xq.data(), 0, yq.data(), 0, 1);
+
+        const float err = fabsf(result - (float)ref);
+        const float tol = 1e-4f * (float)n;
+        if (!(err <= tol)) {
+            printf("FAILED: vec_dot rep %d: got %f, expected %f (err %f)\n", rep, result, (float)ref, err);
+            num_failed++;
+        }
+    }
+
+    return num_failed;
+}
+
+// the scalar reference implementation. The symbol only exists on builds with
+// a native DT3 kernel: without one, arch-fallback.h renames the generic to
+// ggml_vec_dot_dt3_q8_0 and there is nothing to compare against, so the
+// reference is declared weak and the parity test skips when it is absent.
+#if defined(_MSC_VER)
+#define DT3_NO_WEAK_SYMBOLS
+#else
+extern "C" void ggml_vec_dot_dt3_q8_0_generic(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) __attribute__((weak));
+#endif
+
+// the dispatched (possibly vectorized) vec_dot must match the generic scalar
+// implementation exactly — the actual function is called, not a re-derivation
+// of it. Blocks exercise all three regions, the 79/80 and 119/120 boundaries,
+// non-trivial qh bytes (would expose reading their padding 5th digit), and
+// scales of both and mixed signs.
+static int test_vec_dot_arch_parity(const ggml_type_traits_cpu * qfns_cpu) {
+    int num_failed = 0;
+
+#if defined(DT3_NO_WEAK_SYMBOLS)
+    (void) qfns_cpu;
+    printf("(skipping vec_dot arch parity: no weak symbol support)\n");
+    return num_failed;
+#else
+    if (ggml_vec_dot_dt3_q8_0_generic == nullptr) {
+        printf("(skipping vec_dot arch parity: this build has no separate generic vec_dot)\n");
+        return num_failed;
+    }
+
+    const auto * vdot_traits = ggml_get_type_traits_cpu(qfns_cpu->vec_dot_type);
+
+    const int nblocks = 3;
+    const int n       = nblocks*QK_DT3;
+
+    const float scale_cases[][2] = {
+        { 1.0f,   0.25f   },
+        { 0.5f,  -0.125f  }, // negative second plane
+        {-2.0f,   0.75f   }, // negative first plane
+        {-0.75f, -0.0625f }, // both negative
+        { 0.0f,   1.0f    }, // dead first plane
+    };
+    const int n_scale_cases = (int)(sizeof(scale_cases)/sizeof(scale_cases[0]));
+
+    // pattern 0: fully random trits
+    // pattern 1: zero everywhere except elements 120..127 (qh-only)
+    // pattern 2: single +1/-1 walking over the region boundaries
+    const int boundary_pos[] = { 0, 79, 80, 119, 120, 127 };
+
+    for (int rep = 0; rep < 96; ++rep) {
+        std::vector<uint8_t> xq(nblocks*DT3_BLOCK_SIZE);
+
+        for (int i = 0; i < nblocks; ++i) {
+            int8_t t1[QK_DT3] = {0};
+            int8_t t2[QK_DT3] = {0};
+
+            const int pattern = rep % 3;
+            if (pattern == 0) {
+                for (int j = 0; j < QK_DT3; ++j) {
+                    t1[j] = rng_trit();
+                    t2[j] = rng_trit();
+                }
+            } else if (pattern == 1) {
+                for (int j = 120; j < QK_DT3; ++j) {
+                    t1[j] = rng_trit();
+                    t2[j] = rng_trit();
+                }
+            } else {
+                const int pos = boundary_pos[rep/3 % 6];
+                t1[pos] = (rep & 1) ? 1 : -1;
+                t2[QK_DT3 - 1 - pos] = (rep & 1) ? -1 : 1;
+            }
+
+            const float * sc = scale_cases[(rep + i) % n_scale_cases];
+            ref_pack_block(t1, sc[0], t2, sc[1], &xq[i*DT3_BLOCK_SIZE]);
+        }
+
+        std::vector<float> y(n);
+        for (int j = 0; j < n; ++j) {
+            // reach the full q8_0 range, both signs
+            y[j] = 127.0f*sinf(0.7f*(float)(j + 13*rep)) + 0.5f*cosf((float)j);
+        }
+        std::vector<uint8_t> yq(ggml_row_size(qfns_cpu->vec_dot_type, n));
+        vdot_traits->from_float(y.data(), yq.data(), n);
+
+        float res_arch    = INFINITY;
+        float res_generic = -INFINITY;
+        qfns_cpu->vec_dot(n, &res_arch, 0, xq.data(), 0, yq.data(), 0, 1);
+        ggml_vec_dot_dt3_q8_0_generic(n, &res_generic, 0, xq.data(), 0, yq.data(), 0, 1);
+
+        if (memcmp(&res_arch, &res_generic, sizeof(float)) != 0) {
+            printf("FAILED: vec_dot arch parity rep %d: arch %.9g != generic %.9g\n",
+                   rep, res_arch, res_generic);
+            num_failed++;
+        }
+    }
+
+    return num_failed;
+#endif
+}
+
+// --dequant IN.bin OUT.f32 : dequantize raw DT3 blocks, for parity checks
+// against external packers (ternaria's Rust pack_dt3)
+static int run_dequant_file(const char * in_path, const char * out_path) {
+    FILE * fin = fopen(in_path, "rb");
+    if (!fin) {
+        fprintf(stderr, "error: cannot open %s\n", in_path);
+        return 1;
+    }
+    fseek(fin, 0, SEEK_END);
+    const long size = ftell(fin);
+    fseek(fin, 0, SEEK_SET);
+    if (size <= 0 || size % DT3_BLOCK_SIZE != 0) {
+        fprintf(stderr, "error: %s size %ld is not a multiple of %zu\n", in_path, size, DT3_BLOCK_SIZE);
+        fclose(fin);
+        return 1;
+    }
+    std::vector<uint8_t> data(size);
+    if (fread(data.data(), 1, size, fin) != (size_t)size) {
+        fprintf(stderr, "error: short read on %s\n", in_path);
+        fclose(fin);
+        return 1;
+    }
+    fclose(fin);
+
+    const int64_t nel = (int64_t)(size/DT3_BLOCK_SIZE)*QK_DT3;
+    std::vector<float> out(nel);
+    ggml_get_type_traits(GGML_TYPE_DT3)->to_float(data.data(), out.data(), nel);
+
+    FILE * fout = fopen(out_path, "wb");
+    if (!fout) {
+        fprintf(stderr, "error: cannot open %s\n", out_path);
+        return 1;
+    }
+    fwrite(out.data(), sizeof(float), nel, fout);
+    fclose(fout);
+
+    return 0;
+}
+
+int main(int argc, char * argv[]) {
+    if (argc == 4 && strcmp(argv[1], "--dequant") == 0) {
+        return run_dequant_file(argv[2], argv[3]);
+    }
+    if (argc != 1) {
+        fprintf(stderr, "usage: %s [--dequant IN.bin OUT.f32]\n", argv[0]);
+        return 1;
+    }
+
+    ggml_cpu_init();
+
+    const auto * qfns     = ggml_get_type_traits(GGML_TYPE_DT3);
+    const auto * qfns_cpu = ggml_get_type_traits_cpu(GGML_TYPE_DT3);
+
+    int num_failed = 0;
+
+    num_failed += test_layout_constants();
+    num_failed += test_single_trits(qfns);
+    num_failed += test_byte_positions(qfns, qfns_cpu);
+    num_failed += test_validate();
+    num_failed += test_roundtrip(qfns);
+    num_failed += test_quantize_pack_parity(qfns_cpu);
+    num_failed += test_vec_dot(qfns_cpu);
+    num_failed += test_vec_dot_arch_parity(qfns_cpu);
+
+    printf("%d tests failed\n", num_failed);
+
+    return num_failed > 0;
+}

tests/test-quantize-fns.cpp

@@ -158,6 +158,7 @@ static int test_vec_dot_q(bool verbose) {
                 type == GGML_TYPE_Q1_0    ? MAX_QUANTIZATION_TOTAL_ERROR_BINARY :
                 type == GGML_TYPE_TQ1_0   ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY :
                 type == GGML_TYPE_TQ2_0   ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY :
+                type == GGML_TYPE_DT3     ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY :
                 type == GGML_TYPE_Q2_0    ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY :
                 type == GGML_TYPE_Q2_K    ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS :
                 type == GGML_TYPE_IQ2_S   ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS :
@@ -184,7 +185,7 @@ static int test_vec_dot_q(bool verbose) {
                 ? MAX_DOT_PRODUCT_ERROR_LOWBIT
                 : type == GGML_TYPE_Q1_0
                 ? MAX_DOT_PRODUCT_ERROR_BINARY
-                : type == GGML_TYPE_TQ1_0 || type == GGML_TYPE_TQ2_0 || type == GGML_TYPE_Q2_0
+                : type == GGML_TYPE_TQ1_0 || type == GGML_TYPE_TQ2_0 || type == GGML_TYPE_Q2_0 || type == GGML_TYPE_DT3
                 ? MAX_DOT_PRODUCT_ERROR_TERNARY
                 : type == GGML_TYPE_NVFP4
                 ? MAX_DOT_PRODUCT_ERROR_FP4

tools/quantize/quantize.cpp

@@ -47,6 +47,7 @@ static const std::vector<quant_option> QUANT_OPTIONS = {
     { "IQ1_M",    LLAMA_FTYPE_MOSTLY_IQ1_M,    " 1.75 bpw quantization",            },
     { "TQ1_0",    LLAMA_FTYPE_MOSTLY_TQ1_0,    " 1.69 bpw ternarization",           },
     { "TQ2_0",    LLAMA_FTYPE_MOSTLY_TQ2_0,    " 2.06 bpw ternarization",           },
+    { "DT3",      LLAMA_FTYPE_MOSTLY_DT3,      " 3.5 bpw dual-ternary (REF quantizer only - use ternaria's PTQTP for quality)", },
     { "Q2_K",     LLAMA_FTYPE_MOSTLY_Q2_K,     " 2.96G, +3.5199 ppl @ Llama-3-8B",  },
     { "Q2_K_S",   LLAMA_FTYPE_MOSTLY_Q2_K_S,   " 2.96G, +3.1836 ppl @ Llama-3-8B",  },
     { "IQ3_XXS",  LLAMA_FTYPE_MOSTLY_IQ3_XXS,  " 3.06 bpw quantization",            },