cuda : gate DT3 MMQ by batch size, cuBLAS keeps the large-batch prefill

Two integer dot products per weight cancel the 2x int8-over-fp16 tensor
core advantage, so at large batch MMQ cannot beat dequantize + fp16
cuBLAS (measured 331 vs 426 t/s pp512 on the 27B, RTX 4060 Ti, while
running at the same ~20% of its int8 ceiling as Q2_K MMQ does of its
own). MMQ still avoids the dequantization round-trip at moderate batch;
the threshold default is provisional until the crossover is measured
(GGML_CUDA_DT3_MMQ_MAX_BATCH overrides it for that measurement).

ggml/src/ggml-cuda/mmq-config-ampere.cuh

@@ -33,6 +33,23 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
     CASE(GGML_TYPE_Q2_0, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
     CASE(GGML_TYPE_Q2_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
 
+    CASE(GGML_TYPE_DT3, 256, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_DT3, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  24, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  40, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  80, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_DT3, 256, 1, 128,  96, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_DT3, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_DT3, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_DT3, MMQ_ITER_K, true, false);
+
     CASE(GGML_TYPE_Q4_0, 256, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
     CASE(GGML_TYPE_Q4_0, 256, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
     CASE(GGML_TYPE_Q4_0, 256, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);

ggml/src/ggml-cuda/mmq-load-tiles.cuh

@@ -176,6 +176,117 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
     }
 }
 
+template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_dt3(
+        const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
+    constexpr int warp_size   = ggml_cuda_get_physical_warp_size();
+    constexpr int nwarps      = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size;
+    constexpr int I           = ggml_cuda_mmq_get_I(type, J, fallback);
+    constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
+
+    // DT3: 128 elements as two ternary planes with one fp16 scale each, w = d1*t1 + d2*t2.
+    // The two planes cannot be fused into a single int8 value because d1 != d2, so the
+    // tile holds both planes decoded to int8 trits in {-1, 0, +1}, plane 2 offset by
+    // 2*MMQ_TILE_NE_K ints from plane 1 within each row, and 2x8 float scales per row.
+    // The decode is the same base-3 digit iteration as vec_dot_dt3_q8_1 (see vecdotq.cuh):
+    // each packed byte is decoded once with q -> (q*3) & 0xFF, two bytes at a time in the
+    // 16-bit lanes of one int. Digit bytes in {0, 1, 2} become trit bytes in {-1, 0, +1}
+    // without cross-byte borrows via ((dig | 0x80808080) - 0x01010101) ^ 0x80808080.
+#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+    constexpr int row_stride_qs = sram_stride;
+    constexpr int row_stride_df = sram_stride;
+    int   * x_qs = (int   *)  x_tile;
+    float * x_df = (float *) (x_qs + 4*MMQ_TILE_NE_K);
+#else
+    constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_DT3, I);
+    constexpr int row_stride_qs = 4*MMQ_TILE_NE_K + 1;
+    constexpr int row_stride_df = 4*MMQ_TILE_NE_K/QI8_0 + 1;
+    int   * x_qs = (int   *)  x_tile;
+    float * x_df = (float *) (x_qs + txs.qs);
+#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+
+    constexpr int blocks_per_iter = MMQ_ITER_K / QK_DT3;
+    static_assert(blocks_per_iter == 2, "DT3 load assumes 2 blocks per iteration");
+
+    // 32 threads per row: 2 blocks x 2 planes x 8 slots. Slots 0..5 decode one 4-byte
+    // quad of qs each (5 ints of 4 trits), slot 6 decodes the 2 qh bytes (2 ints),
+    // slot 7 is idle.
+    constexpr int threads_per_row = 32;
+    constexpr int nrows = warp_size / threads_per_row;
+
+    const int txi = threadIdx.x % threads_per_row;
+    const int kbx = txi / 16;
+    const int p   = (txi / 8) % 2;
+    const int u   = txi % 8;
+
+#pragma unroll
+    for (int i0 = 0; i0 < I; i0 += nrows*nwarps) {
+        int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row;
+
+        if (fallback) {
+            i = min(i, i_max);
+        }
+
+        const block_dt3 * bxi = (const block_dt3 *) x + kbx0 + i*stride + kbx;
+        int * dst = x_qs + i*row_stride_qs + p*(2*MMQ_TILE_NE_K) + kbx*(QK_DT3/4);
+
+        if (u < 6) {
+            // qs[0..16): byte m holds elements m + 16*n; qs[16..24): byte 16 + m
+            // holds elements 80 + m + 8*n. Either way a quad of consecutive bytes
+            // yields 4 consecutive elements per digit n, i.e. one tile int.
+            const int q32 = get_int_b4(bxi->qs[p], u);
+            int qa = (q32 >> 0) & 0x00FF00FF;
+            int qb = (q32 >> 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 idx = u < 4 ? 4*n + u : 20 + 2*n + (u - 4);
+                dst[idx] = ((dig | 0x80808080) - 0x01010101) ^ 0x80808080;
+            }
+        } else if (u == 6) {
+            // qh: byte b holds elements 120 + b + 2*n for n = 0..3 — only 4 digits
+            // are iterated, the 5th is packing padding and must never decode.
+            int q = bxi->qh[p][0] | (bxi->qh[p][1] << 16);
+#pragma unroll
+            for (int s = 0; s < 2; ++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;
+                }
+                dst[30 + s] = ((dig | 0x80808080) - 0x01010101) ^ 0x80808080;
+            }
+        }
+    }
+
+    // 16 scale entries per row and iteration: 2 planes x 2 blocks x 4 q8_1 chunks.
+    // Plane 2 scales sit after the 8 plane-1 entries, matching the vec_dot indexing.
+    constexpr int scale_entries_per_plane = blocks_per_iter*(QK_DT3/QK8_1);
+
+    const int ksx = threadIdx.x % (2*scale_entries_per_plane);
+    const int ps  = ksx / scale_entries_per_plane;
+    const int scale_block = (ksx % scale_entries_per_plane) / (QK_DT3/QK8_1);
+
+#pragma unroll
+    for (int i0 = 0; i0 < I; i0 += nwarps) {
+        int i = i0 + threadIdx.y;
+
+        if (fallback) {
+            i = min(i, i_max);
+        }
+
+        const block_dt3 * bxi = (const block_dt3 *) x + kbx0 + i*stride + scale_block;
+
+        x_df[i*row_stride_df + ksx] = bxi->d[ps];
+    }
+}
+
 template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_0(
         const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
     constexpr int warp_size   = ggml_cuda_get_physical_warp_size();

ggml/src/ggml-cuda/mmq-vec-dot.cuh

@@ -280,6 +280,140 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma(
 }
 
 
+// DT3: both decoded ternary planes of the tile are multiplied against the same y data,
+// each with its own per-chunk scale: sum = dB * (sumi1*dA1 + sumi2*dA2).
+template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_dt3_q8_1_dp4a(
+        const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) {
+    constexpr int warp_size = ggml_cuda_get_physical_warp_size();
+    constexpr int nwarps    = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size;
+    constexpr int I         = ggml_cuda_mmq_get_I(type, J, fallback);
+
+    constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_DT3, I);
+    const int   * x_qs = (const int   *) x;
+    const float * x_df = (const float *) x_qs + txs.qs;
+    const int   * y_qs = (const int   *) y + 4;
+    const float * y_df = (const float *) y;
+
+    constexpr int row_stride_qs = 4*MMQ_TILE_NE_K + 1;
+    constexpr int row_stride_df = 4*MMQ_TILE_NE_K/QI8_0 + 1;
+
+// #pragma unroll
+    for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += VDR_Q8_0_Q8_1_MMQ) {
+        const int k0 = k00 + k01;
+
+#pragma unroll
+        for (int j0 = 0; j0 < J; j0 += nwarps) {
+            const int j = j0 + threadIdx.y;
+
+#pragma unroll
+            for (int i0 = 0; i0 < I; i0 += warp_size) {
+                const int i = i0 + threadIdx.x;
+
+                const int   * yqs = &y_qs[j*MMQ_TILE_Y_K + k0 % MMQ_TILE_NE_K];
+                const float   dB  =  y_df[j*MMQ_TILE_Y_K + (k0/QI8_1) % (MMQ_TILE_NE_K/QI8_1)];
+
+                sum[j0/nwarps*I/warp_size + i0/warp_size] += vec_dot_q8_0_q8_1_impl<float, VDR_Q8_0_Q8_1_MMQ>
+                    (&x_qs[i*row_stride_qs + k0], yqs,
+                     x_df[i*row_stride_df + k0/QI8_0], dB);
+                sum[j0/nwarps*I/warp_size + i0/warp_size] += vec_dot_q8_0_q8_1_impl<float, VDR_Q8_0_Q8_1_MMQ>
+                    (&x_qs[i*row_stride_qs + 2*MMQ_TILE_NE_K + k0], yqs,
+                     x_df[i*row_stride_df + 2*MMQ_TILE_NE_K/QI8_0 + k0/QI8_0], dB);
+            }
+        }
+    }
+}
+
+template <ggml_type type, int J, bool fallback>
+static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_dt3_q8_1_mma(
+        const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) {
+#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+    // DT3 MMQ is not enabled for AMD (no config entries select it); this only has to compile.
+    GGML_UNUSED_VARS(x, y, sum, k00);
+    NO_DEVICE_CODE;
+#else
+    typedef tile<16, 8, int> tile_A;
+    typedef tile< 8, 8, int> tile_B;
+    typedef tile<16, 8, int> tile_C;
+
+    constexpr int I             = ggml_cuda_mmq_get_I(type, J, fallback);
+    constexpr int sram_stride   = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
+    constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
+    constexpr int ntx           = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
+
+    y += (threadIdx.y % ntx) * (tile_C::J*MMQ_TILE_Y_K);
+
+    const int   * x_qs = (const int   *) x;
+    const float * x_df = (const float *) x_qs + 4*MMQ_TILE_NE_K;
+    const int   * y_qs = (const int   *) y + 4;
+    const float * y_df = (const float *) y;
+
+    tile_A A[ntx][2][MMQ_TILE_NE_K/QI8_0];
+    float dA[ntx][tile_C::ne/2][2][MMQ_TILE_NE_K/QI8_0];
+
+    const int i0 = (threadIdx.y/ntx)*rows_per_warp;
+
+#pragma unroll
+    for (int n = 0; n < ntx; ++n) {
+#pragma unroll
+        for (int p = 0; p < 2; ++p) {
+#pragma unroll
+            for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += QI8_0) {
+                const int k0 = k00 + k01;
+
+                load_ldmatrix(A[n][p][k01/QI8_0], x_qs + (i0 + n*tile_A::I)*sram_stride + p*(2*MMQ_TILE_NE_K) + k0, sram_stride);
+            }
+        }
+
+#pragma unroll
+        for (int l = 0; l < tile_C::ne/2; ++l) {
+            const int i = i0 + n*tile_A::I + tile_C::get_i(2*l);
+
+#pragma unroll
+            for (int p = 0; p < 2; ++p) {
+#pragma unroll
+                for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += QI8_0) {
+                    const int k0 = k00 + k01;
+
+                    dA[n][l][p][k01/QI8_0] = x_df[i*sram_stride + p*(2*MMQ_TILE_NE_K/QI8_0) + k0/QI8_0];
+                }
+            }
+        }
+    }
+
+#pragma unroll
+    for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) {
+#pragma unroll
+        for (int k01 = 0; k01 < MMQ_TILE_NE_K; k01 += QI8_0) {
+            tile_B B;
+            float dB[tile_C::ne/2];
+
+            load_generic(B, y_qs + j0*MMQ_TILE_Y_K + k01, MMQ_TILE_Y_K); // faster than load_ldmatrix
+
+#pragma unroll
+            for (int l = 0; l < tile_C::ne/2; ++l) {
+                const int j = j0 + tile_C::get_j(l);
+
+                dB[l] = y_df[j*MMQ_TILE_Y_K + k01/QI8_1];
+            }
+
+#pragma unroll
+            for (int n = 0; n < ntx; ++n) {
+                tile_C C1;
+                tile_C C2;
+                mma(C1, A[n][0][k01/QI8_0], B);
+                mma(C2, A[n][1][k01/QI8_0], B);
+
+#pragma unroll
+                for (int l = 0; l < tile_C::ne; ++l) {
+                    sum[(j0/tile_C::J + n)*tile_C::ne + l] +=
+                        (C1.x[l]*dA[n][l/2][0][k01/QI8_0] + C2.x[l]*dA[n][l/2][1][k01/QI8_0])*dB[l%2];
+                }
+            }
+        }
+    }
+#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+}
+
 template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_1_q8_1_dp4a(
         const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) {
     constexpr int warp_size = ggml_cuda_get_physical_warp_size();

ggml/src/ggml-cuda/mmq.cu

@@ -4,6 +4,7 @@
 #include "mmid.cuh"
 
 #include <cstdint>
+#include <cstdlib>
 
 static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) {
     switch (args.type_x) {
@@ -13,6 +14,9 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con
         case GGML_TYPE_Q2_0:
             mul_mat_q_case<GGML_TYPE_Q2_0>(ctx, args, stream);
             break;
+        case GGML_TYPE_DT3:
+            mul_mat_q_case<GGML_TYPE_DT3>(ctx, args, stream);
+            break;
         case GGML_TYPE_Q4_0:
             mul_mat_q_case<GGML_TYPE_Q4_0>(ctx, args, stream);
             break;
@@ -261,6 +265,30 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
     return false;
 #endif // GGML_CUDA_FORCE_CUBLAS
 
+    // DT3 keeps two decoded ternary planes per row in SRAM, roughly double the tile
+    // of a single-plane type: only the MMA data layout is implemented and even the
+    // narrowest tile needs ~76 KiB of shared memory, more than e.g. Turing offers.
+    if (type == GGML_TYPE_DT3) {
+        if (!turing_mma_available(cc)) {
+            return false;
+        }
+        // Two integer dot products per weight cancel the 2x int8-over-fp16 advantage
+        // of the tensor cores, so at large batch the dequantize + fp16 cuBLAS path
+        // wins; MMQ avoids the dequantization round-trip and wins below the
+        // crossover (measured on RTX 4060 Ti). Override for experiments with
+        // GGML_CUDA_DT3_MMQ_MAX_BATCH.
+        static const int64_t max_batch = []() {
+            const char * env = getenv("GGML_CUDA_DT3_MMQ_MAX_BATCH");
+            return env ? atoll(env) : 192;
+        }();
+        if (ne11 > max_batch) {
+            return false;
+        }
+        const int    id    = ggml_cuda_get_device();
+        const size_t smpbo = ggml_cuda_info().devices[id].smpbo;
+        return mmq_get_nbytes_shared(ggml_cuda_mmq_get_config(GGML_TYPE_DT3, 8, true, cc), cc) <= smpbo;
+    }
+
     bool mmq_supported;
 
     switch (type) {

ggml/src/ggml-cuda/mmq.cuh

@@ -61,6 +61,7 @@ static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) {
     switch (type_x) {
         case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q2_0:
+        case GGML_TYPE_DT3:
             return MMQ_Q8_1_DS_LAYOUT_D4;
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:
@@ -121,6 +122,7 @@ struct tile_x_sizes {
 
 enum ggml_cuda_mmq_sram_layout {
     GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0,
+    GGML_CUDA_MMQ_SRAM_LAYOUT_DT3,   // Two decoded ternary planes per row, each with its own per-block scales.
     GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1,
     GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K,
     GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K,
@@ -133,6 +135,8 @@ static constexpr __host__ __device__ int ggml_cuda_mmq_get_sram_stride(ggml_cuda
     switch (sram_layout) {
         case GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0:
             return 2*MMQ_TILE_NE_K + 2*MMQ_TILE_NE_K/QI8_0 + 4;
+        case GGML_CUDA_MMQ_SRAM_LAYOUT_DT3:
+            return 4*MMQ_TILE_NE_K + 4*MMQ_TILE_NE_K/QI8_0 + 4;
         case GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1:
             return 2*MMQ_TILE_NE_K + 2*MMQ_TILE_NE_K/QI8_1 + 4;
         case GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K:
@@ -151,6 +155,7 @@ static constexpr __host__ __device__ int ggml_cuda_mmq_get_sram_stride(ggml_cuda
 }
 
 static_assert(ggml_cuda_mmq_get_sram_stride(GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0)  % 8 == 4, "Wrong padding.");
+static_assert(ggml_cuda_mmq_get_sram_stride(GGML_CUDA_MMQ_SRAM_LAYOUT_DT3)   % 8 == 4, "Wrong padding.");
 static_assert(ggml_cuda_mmq_get_sram_stride(GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1)  % 8 == 4, "Wrong padding.");
 static_assert(ggml_cuda_mmq_get_sram_stride(GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K)  % 8 == 4, "Wrong padding.");
 static_assert(ggml_cuda_mmq_get_sram_stride(GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K)  % 8 == 4, "Wrong padding.");
@@ -377,6 +382,7 @@ static constexpr __device__ int ggml_cuda_mmq_get_rows_per_warp(ggml_type type,
 #define MMQ_DP4A_TXS_Q8_0    tile_x_sizes{I*MMQ_TILE_NE_K*2 + I, I*MMQ_TILE_NE_K*2/QI8_0 + I/(QI8_0/2), 0}
 #define MMQ_DP4A_TXS_Q8_0_16 tile_x_sizes{I*MMQ_TILE_NE_K*2 + I, I*MMQ_TILE_NE_K*4/QI8_0 + I/(QI8_0/4), 0}
 #define MMQ_DP4A_TXS_Q8_1    tile_x_sizes{I*MMQ_TILE_NE_K*2 + I, I*MMQ_TILE_NE_K*2/QI8_1 + I/(QI8_1/2), 0}
+#define MMQ_DP4A_TXS_DT3     tile_x_sizes{I*MMQ_TILE_NE_K*4 + I, I*MMQ_TILE_NE_K*4/QI8_0 + I,           0}
 #define MMQ_DP4A_TXS_Q2_K    tile_x_sizes{I*MMQ_TILE_NE_K*2 + I, I*MMQ_TILE_NE_K         + I,           0}
 #define MMQ_DP4A_TXS_Q3_K    tile_x_sizes{I*MMQ_TILE_NE_K*2 + I, I,                                     I*MMQ_TILE_NE_K/8 + I/8}
 #define MMQ_DP4A_TXS_Q4_K    tile_x_sizes{I*MMQ_TILE_NE_K   + I, I*MMQ_TILE_NE_K/QI4_K,                 I*MMQ_TILE_NE_K/8 + I/8}
@@ -387,6 +393,7 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml
     switch (type) {
         case GGML_TYPE_Q1_0:    return MMQ_DP4A_TXS_Q8_0;
         case GGML_TYPE_Q2_0:    return MMQ_DP4A_TXS_Q8_0;
+        case GGML_TYPE_DT3:     return MMQ_DP4A_TXS_DT3;
         case GGML_TYPE_Q4_0:    return MMQ_DP4A_TXS_Q4_0;
         case GGML_TYPE_Q4_1:    return MMQ_DP4A_TXS_Q4_1;
         case GGML_TYPE_Q5_0:    return MMQ_DP4A_TXS_Q8_0;
@@ -550,6 +557,12 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func
                     ggml_cuda_mmq_load_tiles_q2_0<type, J, fallback>,
                     ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a<type, J, fallback>,
                     ggml_cuda_mmq_write_back_dp4a<type, J, fallback>);
+            case GGML_TYPE_DT3:
+                return ggml_cuda_mmq_util_funcs(
+                    VDR_Q8_0_Q8_1_MMQ,
+                    ggml_cuda_mmq_load_tiles_dt3<type, J, fallback>,
+                    ggml_cuda_mmq_vec_dot_dt3_q8_1_dp4a<type, J, fallback>,
+                    ggml_cuda_mmq_write_back_dp4a<type, J, fallback>);
             case GGML_TYPE_Q4_0:
                 return ggml_cuda_mmq_util_funcs(
                     VDR_Q4_0_Q8_1_MMQ,
@@ -714,6 +727,12 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func
                 ggml_cuda_mmq_load_tiles_q2_0<type, J, fallback>,
                 ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma<type, J, fallback, MMQ_Q8_1_DS_LAYOUT_D4>,
                 ggml_cuda_mmq_write_back_mma<type, J, fallback>);
+        case GGML_TYPE_DT3:
+            return ggml_cuda_mmq_util_funcs(
+                -1,
+                ggml_cuda_mmq_load_tiles_dt3<type, J, fallback>,
+                ggml_cuda_mmq_vec_dot_dt3_q8_1_mma<type, J, fallback>,
+                ggml_cuda_mmq_write_back_mma<type, J, fallback>);
         case GGML_TYPE_Q4_0:
             return ggml_cuda_mmq_util_funcs(
                 -1,
@@ -1565,6 +1584,7 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda
 
 extern DECL_MMQ_CASE(GGML_TYPE_Q1_0);
 extern DECL_MMQ_CASE(GGML_TYPE_Q2_0);
+extern DECL_MMQ_CASE(GGML_TYPE_DT3);
 extern DECL_MMQ_CASE(GGML_TYPE_Q4_0);
 extern DECL_MMQ_CASE(GGML_TYPE_Q4_1);
 extern DECL_MMQ_CASE(GGML_TYPE_Q5_0);

ggml/src/ggml-cuda/template-instances/generate_cu_files.py

@@ -37,6 +37,7 @@ SOURCE_FATTN_MMA_CASE = "DECL_FATTN_MMA_F16_CASE({head_size_kq}, {head_size_v},
 TYPES_MMQ = [
     "GGML_TYPE_Q1_0",
     "GGML_TYPE_Q2_0",
+    "GGML_TYPE_DT3",
     "GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q8_0",
     "GGML_TYPE_Q2_K", "GGML_TYPE_Q3_K", "GGML_TYPE_Q4_K", "GGML_TYPE_Q5_K", "GGML_TYPE_Q6_K",
     "GGML_TYPE_IQ2_XXS", "GGML_TYPE_IQ2_XS", "GGML_TYPE_IQ2_S", "GGML_TYPE_IQ3_XXS", "GGML_TYPE_IQ3_S",

ggml/src/ggml-cuda/template-instances/mmq-instance-dt3.cu

@@ -0,0 +1,5 @@
+// This file has been autogenerated by generate_cu_files.py, do not edit manually.
+
+#include "../mmq.cuh"
+
+DECL_MMQ_CASE(GGML_TYPE_DT3);

ggml/src/ggml-vulkan/ggml-vulkan.cpp

@@ -5220,6 +5220,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_BF16][i], "mul_mat_vec_bf16_f32_f32", arr_dmmv_bf16_f32_f32_len[reduc], arr_dmmv_bf16_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {wg_size_subgroup, 2, i+1}, 1, false, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q1_0][i], "mul_mat_vec_q1_0_f32_f32", arr_dmmv_q1_0_f32_f32_len[reduc], arr_dmmv_q1_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q2_0][i], "mul_mat_vec_q2_0_f32_f32", arr_dmmv_q2_0_f32_f32_len[reduc], arr_dmmv_q2_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
+            ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_DT3 ][i], "mul_mat_vec_dt3_f32_f32",  arr_dmmv_dt3_f32_f32_len[reduc],  arr_dmmv_dt3_f32_f32_data[reduc],  "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q4_0][i], "mul_mat_vec_q4_0_f32_f32", arr_dmmv_q4_0_f32_f32_len[reduc], arr_dmmv_q4_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q4_1][i], "mul_mat_vec_q4_1_f32_f32", arr_dmmv_q4_1_f32_f32_len[reduc], arr_dmmv_q4_1_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q5_0][i], "mul_mat_vec_q5_0_f32_f32", arr_dmmv_q5_0_f32_f32_len[reduc], arr_dmmv_q5_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
@@ -5247,6 +5248,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_BF16][i], "mul_mat_vec_bf16_f16_f32", arr_dmmv_bf16_f16_f32_len[reduc], arr_dmmv_bf16_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {wg_size_subgroup, 2, i+1}, 1, false, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q1_0][i], "mul_mat_vec_q1_0_f16_f32", arr_dmmv_q1_0_f16_f32_len[reduc], arr_dmmv_q1_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q2_0][i], "mul_mat_vec_q2_0_f16_f32", arr_dmmv_q2_0_f16_f32_len[reduc], arr_dmmv_q2_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
+            ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_DT3 ][i], "mul_mat_vec_dt3_f16_f32",  arr_dmmv_dt3_f16_f32_len[reduc],  arr_dmmv_dt3_f16_f32_data[reduc],  "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q4_0][i], "mul_mat_vec_q4_0_f16_f32", arr_dmmv_q4_0_f16_f32_len[reduc], arr_dmmv_q4_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q4_1][i], "mul_mat_vec_q4_1_f16_f32", arr_dmmv_q4_1_f16_f32_len[reduc], arr_dmmv_q4_1_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q5_0][i], "mul_mat_vec_q5_0_f16_f32", arr_dmmv_q5_0_f16_f32_len[reduc], arr_dmmv_q5_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size);
@@ -5362,6 +5364,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_F32 ], "f32_to_f16",   dequant_f32_len,  dequant_f32_data,  "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q1_0], "dequant_q1_0", dequant_q1_0_len, dequant_q1_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 8, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q2_0], "dequant_q2_0", dequant_q2_0_len, dequant_q2_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 4, 1, 1}, {}, 1);
+    ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_DT3 ], "dequant_dt3",  dequant_dt3_len,  dequant_dt3_data,  "main", 2, 5 * sizeof(uint32_t), {256 * 4, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q4_0], "dequant_q4_0", dequant_q4_0_len, dequant_q4_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q4_1], "dequant_q4_1", dequant_q4_1_len, dequant_q4_1_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q5_0], "dequant_q5_0", dequant_q5_0_len, dequant_q5_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
@@ -5390,6 +5393,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_BF16], "get_rows_bf16", get_rows_bf16_len, get_rows_bf16_data, "main", 3, sizeof(vk_op_binary_push_constants), { 512, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q1_0], "get_rows_q1_0", get_rows_q1_0_len, get_rows_q1_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q2_0], "get_rows_q2_0", get_rows_q2_0_len, get_rows_q2_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+    ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_DT3 ], "get_rows_dt3",  get_rows_dt3_len,  get_rows_dt3_data,  "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q4_0], "get_rows_q4_0", get_rows_q4_0_len, get_rows_q4_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q4_1], "get_rows_q4_1", get_rows_q4_1_len, get_rows_q4_1_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q5_0], "get_rows_q5_0", get_rows_q5_0_len, get_rows_q5_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
@@ -5418,6 +5422,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_BF16], "get_rows_bf16_f32", get_rows_bf16_f32_len, get_rows_bf16_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), { 512, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q1_0], "get_rows_q1_0_f32", get_rows_q1_0_f32_len, get_rows_q1_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q2_0], "get_rows_q2_0_f32", get_rows_q2_0_f32_len, get_rows_q2_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+    ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_DT3 ], "get_rows_dt3_f32",  get_rows_dt3_f32_len,  get_rows_dt3_f32_data,  "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q4_0], "get_rows_q4_0_f32", get_rows_q4_0_f32_len, get_rows_q4_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q4_1], "get_rows_q4_1_f32", get_rows_q4_1_f32_len, get_rows_q4_1_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q5_0], "get_rows_q5_0_f32", get_rows_q5_0_f32_len, get_rows_q5_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
@@ -7617,6 +7622,7 @@ static vk_pipeline ggml_vk_get_to_fp16(ggml_backend_vk_context * ctx, ggml_type
         case GGML_TYPE_F32:
         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:
@@ -7760,6 +7766,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec(ggml_backend_vk_context *
         case GGML_TYPE_BF16:
         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:
@@ -9171,12 +9178,18 @@ static void ggml_vk_mul_mat_q_f16(ggml_backend_vk_context * ctx, vk_context& sub
 
     bool quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && !y_non_contig && (ne11 * ne10) % 4 == 0;
 
+    // DT3 weights (d1*t1 + d2*t2, two fp16-scaled ternary planes) already pay
+    // one fp16 rounding in the dequant fallback; fp16 accumulation on top of
+    // it costs measurable perplexity. Force fp32 accumulators, matching the
+    // numerics of the CUDA GEMM fallback (fp16 inputs, fp32 compute).
+    const ggml_prec mm_prec = src0->type == GGML_TYPE_DT3 ? GGML_PREC_F32 : (ggml_prec)dst->op_params[0];
+
     // Check for mmq first
-    vk_matmul_pipeline mmp = quantize_y ? ggml_vk_get_mul_mat_mat_pipeline(ctx, src0->type, GGML_TYPE_Q8_1, (ggml_prec)dst->op_params[0]) : nullptr;
+    vk_matmul_pipeline mmp = quantize_y ? ggml_vk_get_mul_mat_mat_pipeline(ctx, src0->type, GGML_TYPE_Q8_1, mm_prec) : nullptr;
 
     if (mmp == nullptr) {
         // Fall back to f16 dequant mul mat
-        mmp = ggml_vk_get_mul_mat_mat_pipeline(ctx, src0->type, y_non_contig ? f16_type : src1->type, (ggml_prec)dst->op_params[0]);
+        mmp = ggml_vk_get_mul_mat_mat_pipeline(ctx, src0->type, y_non_contig ? f16_type : src1->type, mm_prec);
         quantize_y = false;
     }
 
@@ -9185,7 +9198,7 @@ static void ggml_vk_mul_mat_q_f16(ggml_backend_vk_context * ctx, vk_context& sub
 
     if (qx_needs_dequant) {
         // Fall back to dequant + f16 mulmat
-        mmp = ggml_vk_get_mul_mat_mat_pipeline(ctx, f16_type, y_f32_kernel ? GGML_TYPE_F32 : f16_type, (ggml_prec)dst->op_params[0]);
+        mmp = ggml_vk_get_mul_mat_mat_pipeline(ctx, f16_type, y_f32_kernel ? GGML_TYPE_F32 : f16_type, mm_prec);
     }
 
     // Not implemented
@@ -17987,6 +18000,13 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
                     }
                 }
                 switch (src0_type) {
+                    case GGML_TYPE_DT3:
+                        // DT3 has dequant, get_rows and scalar mul_mat_vec shaders only:
+                        // mul_mat_id, coopmat and MMQ are intentionally not implemented
+                        if (op->op == GGML_OP_MUL_MAT_ID) {
+                            return false;
+                        }
+                        break;
                     case GGML_TYPE_F32:
                     case GGML_TYPE_F16:
                     case GGML_TYPE_BF16:
@@ -18097,6 +18117,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
                     case GGML_TYPE_BF16:
                     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-vulkan/vulkan-shaders/dequant_dt3.comp

@@ -0,0 +1,47 @@
+#version 450
+
+#include "dequant_head.glsl"
+
+layout(local_size_x = 256, local_size_y = 1, local_size_z = 1) in;
+
+layout (binding = 0) readonly buffer A {block_dt3 data_a[];};
+layout (binding = 1) writeonly buffer D {D_TYPE data_b[];};
+
+// Eight blocks per workgroup, 32 threads per block. Threads 0..23 decode one
+// qs byte of each plane (5 trits in base 3), threads 24..25 decode one qh
+// byte of each plane (4 trits — the 5th base-3 digit of a qh byte is packing
+// padding that always decodes to -1, so it must not be read), threads 26..31
+// idle.
+void main() {
+    const uint ib = gl_WorkGroupID.x * 8 + gl_LocalInvocationID.x / 32;
+    const uint il = gl_LocalInvocationID.x % 32;
+
+    if (ib >= p.nel / 128 || il >= 26) {
+        return;
+    }
+
+    const float d1 = float(data_a[ib].d[0]);
+    const float d2 = float(data_a[ib].d[1]);
+
+    const uint b_idx = ib * 128;
+
+    // element covered by the first digit, distance between consecutive
+    // digits, and number of digits stored in this byte
+    const bool is_qh  = il >= 24;
+    const uint e0     = is_qh ? 120 + (il - 24) : (il < 16 ? il : 80 + (il - 16));
+    const uint stride = is_qh ? 2 : (il < 16 ? 16 : 8);
+    const uint digits = is_qh ? 4 : 5;
+
+    uint q1 = is_qh ? uint(data_a[ib].qh[il - 24])     : uint(data_a[ib].qs[il]);
+    uint q2 = is_qh ? uint(data_a[ib].qh[2 + il - 24]) : uint(data_a[ib].qs[24 + il]);
+
+    // decode each byte once: take the top base-3 digit with (q*3) >> 8, then
+    // shift it out with q <- (q*3) mod 256
+    for (uint n = 0; n < digits; ++n) {
+        const float t1 = float(int((q1 * 3) >> 8) - 1);
+        const float t2 = float(int((q2 * 3) >> 8) - 1);
+        data_b[b_idx + e0 + n*stride] = D_TYPE(d1*t1 + d2*t2);
+        q1 = (q1 * 3) & 0xFF;
+        q2 = (q2 * 3) & 0xFF;
+    }
+}

ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl

@@ -154,6 +154,49 @@ vec4 dequantize4(uint ib, uint iqs, uint a_offset) {
 }
 #endif
 
+#if defined(DATA_A_DT3)
+// Dual-plane ternary: element iqs of plane p sits in a base-3 packed byte.
+// Elements 0..79 use qs[m], m = iqs % 16, digit n = iqs / 16; elements
+// 80..119 use qs[16 + m], m = (iqs - 80) % 8, digit n = (iqs - 80) / 8;
+// elements 120..127 use qh[j], j = iqs % 2, digit n = (iqs - 120) / 2.
+// A qh byte holds only 4 trits: its 5th base-3 digit is packing padding that
+// always decodes to -1, never to 0, so it must not be read.
+// The decode multiplies the byte by 3^n modulo 256 and takes the top digit.
+float dt3_get_trit(uint ib, uint p, uint iqs, uint a_offset) {
+    const uint pow3[5] = {1, 3, 9, 27, 81};
+    uint b;
+    uint n;
+    if (iqs < 80) {
+        b = uint(data_a[a_offset + ib].qs[p*24 + (iqs & 15)]);
+        n = iqs >> 4;
+    } else if (iqs < 120) {
+        b = uint(data_a[a_offset + ib].qs[p*24 + 16 + ((iqs - 80) & 7)]);
+        n = (iqs - 80) >> 3;
+    } else {
+        b = uint(data_a[a_offset + ib].qh[p*2 + (iqs & 1)]);
+        n = (iqs - 120) >> 1;
+    }
+    const uint q = (b * pow3[n]) & 0xFF;
+    return float(int((q * 3) >> 8) - 1);
+}
+// w = d1*t1 + d2*t2; both products are exact (t in {-1,0,+1}), so the sum has
+// a single float rounding and matches the CPU reference bit by bit
+vec2 dequantize(uint ib, uint iqs, uint a_offset) {
+    const float d1 = float(data_a[a_offset + ib].d[0]);
+    const float d2 = float(data_a[a_offset + ib].d[1]);
+    return vec2(d1*dt3_get_trit(ib, 0, iqs,     a_offset) + d2*dt3_get_trit(ib, 1, iqs,     a_offset),
+                d1*dt3_get_trit(ib, 0, iqs + 1, a_offset) + d2*dt3_get_trit(ib, 1, iqs + 1, a_offset));
+}
+vec4 dequantize4(uint ib, uint iqs, uint a_offset) {
+    const float d1 = float(data_a[a_offset + ib].d[0]);
+    const float d2 = float(data_a[a_offset + ib].d[1]);
+    return vec4(d1*dt3_get_trit(ib, 0, iqs,     a_offset) + d2*dt3_get_trit(ib, 1, iqs,     a_offset),
+                d1*dt3_get_trit(ib, 0, iqs + 1, a_offset) + d2*dt3_get_trit(ib, 1, iqs + 1, a_offset),
+                d1*dt3_get_trit(ib, 0, iqs + 2, a_offset) + d2*dt3_get_trit(ib, 1, iqs + 2, a_offset),
+                d1*dt3_get_trit(ib, 0, iqs + 3, a_offset) + d2*dt3_get_trit(ib, 1, iqs + 3, a_offset));
+}
+#endif
+
 #if defined(DATA_A_IQ1_S)
 vec2 dequantize(uint ib, uint iqs, uint a_offset) {
     const uint ib32 = iqs / 32;
@@ -571,6 +614,13 @@ vec2 get_dm(uint ib, uint a_offset) {
 }
 #endif
 
+#if defined(DATA_A_DT3)
+// the two scales are already applied inside dequantize/dequantize4
+vec2 get_dm(uint ib, uint a_offset) {
+    return vec2(1, 0);
+}
+#endif
+
 #if defined(DATA_A_MXFP4)
 vec2 get_dm(uint ib, uint a_offset) {
     return vec2(e8m0_to_fp32(data_a[a_offset + ib].e), 0);

ggml/src/ggml-vulkan/vulkan-shaders/types.glsl

@@ -235,6 +235,27 @@ struct block_q2_0_packed16
 #define DATA_A_QUANT_LEGACY
 #endif
 
+#define QUANT_K_DT3 128
+#define QUANT_R_DT3 1
+
+// Dual-plane ternary: w = d[0]*t1 + d[1]*t2 with trits in {-1,0,+1}.
+// Per plane: 24 bytes with 5 trits each in base 3 (elements 0..119), then
+// 2 bytes with 4 trits each (elements 120..127). Plane p uses qs[p*24..],
+// qh[p*2..] and d[p].
+struct block_dt3
+{
+    uint8_t qs[2*24];
+    uint8_t qh[2*2];
+    float16_t d[2];
+};
+
+#if defined(DATA_A_DT3)
+#define QUANT_K QUANT_K_DT3
+#define QUANT_R QUANT_R_DT3
+#define QUANT_AUXF 1
+#define A_TYPE block_dt3
+#endif
+
 #define QUANT_K_Q8_1 32
 #define QUANT_R_Q8_1 1
 

ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp

@@ -51,6 +51,7 @@ const std::vector<std::string> type_names = {
     "f16",
     "q1_0",
     "q2_0",
+    "dt3",
     "q4_0",
     "q4_1",
     "q5_0",
@@ -591,6 +592,12 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c
             continue;
         }
 
+        // DT3 has no direct matmul shaders: mul_mat goes through dequant to
+        // f16 + f16 matmul, and coopmat/MMQ are intentionally not implemented
+        if (tname == "dt3") {
+            continue;
+        }
+
         std::string data_a_key = "DATA_A_" + to_uppercase(tname);
         // For aligned matmul loads
         std::string load_vec_a = (coopmat2 || tname == "f32" || tname == "f16" || tname == "bf16") ? load_vec : load_vec_quant;
@@ -758,9 +765,12 @@ void process_shaders() {
         }
 #endif
 
-        string_to_spv("mul_mat_vec_id_" + tname + "_f32_f32", shader, merge_maps(base_dict, {{"MUL_MAT_ID", "1"}, {data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPEV2", "vec2"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}));
-        string_to_spv("mul_mat_vec_id_" + tname + "_f32_f32_subgroup", shader, merge_maps(base_dict, {{"MUL_MAT_ID", "1"}, {data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPEV2", "vec2"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}, {"USE_SUBGROUP_ADD", "1"}}));
-        string_to_spv("mul_mat_vec_id_" + tname + "_f32_f32_subgroup_no_shmem", shader, merge_maps(base_dict, {{"MUL_MAT_ID", "1"}, {data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPEV2", "vec2"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}, {"USE_SUBGROUP_ADD_NO_SHMEM", "1"}}));
+        // mul_mat_id is not implemented for DT3 (supports_op declines it)
+        if (tname != "dt3") {
+            string_to_spv("mul_mat_vec_id_" + tname + "_f32_f32", shader, merge_maps(base_dict, {{"MUL_MAT_ID", "1"}, {data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPEV2", "vec2"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}));
+            string_to_spv("mul_mat_vec_id_" + tname + "_f32_f32_subgroup", shader, merge_maps(base_dict, {{"MUL_MAT_ID", "1"}, {data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPEV2", "vec2"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}, {"USE_SUBGROUP_ADD", "1"}}));
+            string_to_spv("mul_mat_vec_id_" + tname + "_f32_f32_subgroup_no_shmem", shader, merge_maps(base_dict, {{"MUL_MAT_ID", "1"}, {data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPEV2", "vec2"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}, {"USE_SUBGROUP_ADD_NO_SHMEM", "1"}}));
+        }
 
         // mul mat vec with integer dot product
 #if defined(GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT)
@@ -1254,7 +1264,8 @@ void write_output_files() {
             src << "const uint64_t arr_dmmv_" << tname << "_" << btype << "_f32_len[3] =  {mul_mat_vec_" << tname << "_" << btype << "_f32_len,  mul_mat_vec_" << tname << "_" << btype << "_f32_subgroup_len, mul_mat_vec_"  << tname << "_" << btype << "_f32_subgroup_no_shmem_len};\n";
         }
 
-        if (btype == "f16") {
+        if (btype == "f16" || tname == "dt3") {
+            // no mul_mat_vec_id shaders for DT3
             continue;
         }
         hdr << "extern const void * arr_dmmv_id_"   << tname << "_" << btype << "_f32_data[3];\n";

tests/test-dt3-gpu.cpp

@@ -18,7 +18,10 @@
 // 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
+// MUL_MAT with more destination columns than the MMVQ limit takes the MMQ
+// path where the backend implements it for DT3 (CUDA on Ampere-class
+// hardware and newer): integer dot products in the same numerical regime as
+// MMVQ, judged just as strictly. Backends without DT3 MMQ fall 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
@@ -33,8 +36,12 @@
 // 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.
+// DT3 is implemented for CUDA, HIP and Vulkan. Without one of those backends
+// the test is skipped and succeeds — an unsupported backend is not a failure.
+// On Vulkan the n <= 8 path is the scalar mul_mat_vec shader (fp32 dot on
+// exactly decoded weights, not an integer dot), and the larger-n path is
+// dequantization to fp16 + the f16 matmul pipeline; both are judged by the
+// same gates as the CUDA MMVQ/GEMM paths.
 
 #include "ggml.h"
 #include "ggml-alloc.h"
@@ -260,7 +267,8 @@ static int test_mul_mat(ggml_backend_t backend, ggml_type type, const std::vecto
                         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};
+    // 100 exercises a wide MMQ tile with a clamped last column block
+    const int ncols_dst[] = {1, 2, 5, 8, 16, 100};
 
     // the same weights as the fp16 GEMM fallback sees them
     std::vector<float> ref_w16(ref_w.size());
@@ -270,6 +278,8 @@ static int test_mul_mat(ggml_backend_t backend, ggml_type type, const std::vecto
 
     std::vector<std::vector<float>> results;
 
+    bool n16_integer = false; // whether the ncols_dst = 16 run took an integer (MMQ) path
+
     for (int c = 0; c < (int)(sizeof(ncols_dst)/sizeof(ncols_dst[0])); ++c) {
         const int n = ncols_dst[c];
 
@@ -320,25 +330,33 @@ static int test_mul_mat(ggml_backend_t backend, ggml_type type, const std::vecto
         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
+        // against the exact reference (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). Larger n takes
+        // the MMQ path where the backend implements it for this type: integer
+        // dot products in the same numerical regime as MMVQ, judged just as
+        // strictly. Backends without MMQ for the type fall back to dequantize
+        // + GEMM, 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);
+        // ours: for the strict type that run is gated below by bit-identity
+        // with the same GEMM on an F16 tensor, and only reported here.
+        // The two regimes are told apart by the result itself: an integer path
+        // lands within float rounding of the exact reference, a fp16/TF32 GEMM
+        // stays orders of magnitude above it. A broken MMQ kernel cannot hide
+        // in the GEMM class: it would then have to be bit-identical to the F16
+        // GEMM control below, which an integer path never is.
+        const bool   is_mmvq  = n <= 8;
+        const bool   integer_path = is_mmvq || err.norm_rel <= 1e-5;
+        if (n == 16) {
+            n16_integer = integer_path;
+        }
+        const bool   gated   = integer_path || !strict;
+        const double err_gate = integer_path ? 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",
+        printf("%s: %s mul_mat GPU, ncols_dst = %3d (%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" : integer_path ? "MMQ" : "GEMM",
                err.norm_rel, err16.norm_rel, err.max_rel);
         if (failed) {
             num_failed++;
@@ -348,24 +366,26 @@ static int test_mul_mat(ggml_backend_t backend, ggml_type type, const std::vecto
         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
+    // MMVQ vs the batched path: the first 8 columns of the ncols_dst = 16 run
+    // must match the ncols_dst = 8 MMVQ run. When the batched run took the
+    // integer MMQ path both sides are exact to float rounding of the
+    // accumulation; against a GEMM fallback the gate is fp16 weight rounding.
     {
-        const std::vector<float> & mmvq = results[3]; // n = 8
-        const std::vector<float> & gemm = results[4]; // n = 16
+        const std::vector<float> & mmvq  = results[3]; // n = 8
+        const std::vector<float> & batch = 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];
+                const double diff = (double)mmvq[(size_t)j*NROWS + r] - (double)batch[(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);
+        const double tol      = !strict ? 1e-2 : n16_integer ? 1e-5 : 5e-3;
+        printf("%s: %s MMVQ vs %s path on shared columns: norm rel err = %g\n",
+               norm_rel <= tol ? "OK" : "FAILED", ggml_type_name(type), n16_integer ? "MMQ" : "GEMM", norm_rel);
         if (norm_rel > tol) {
             num_failed++;
         }
@@ -375,50 +395,73 @@ static int test_mul_mat(ggml_backend_t backend, ggml_type type, const std::vecto
     // 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);
+    // dequantization from cuBLAS numerics. The backend may run the DT3
+    // fallback at a different accumulator precision than its default F16
+    // GEMM (Vulkan forces fp32 accumulators for DT3), so the F16 control is
+    // run at both the default and the F32-forced precision and bit-identity
+    // with either one passes. When the ncols_dst = 16 run took the integer
+    // MMQ path there is no dequantization involved and no GEMM to compare
+    // against — that run was already gated strictly above.
+    if (strict && n16_integer) {
+        printf("OK: %s ncols_dst = 16 took the integer MMQ path, F16 GEMM bit-identity control not applicable\n",
+               ggml_type_name(type));
+    }
+    if (strict && !n16_integer) {
+        int    n_mismatch_best = -1;
+        double max_diff_best   = 0.0;
 
-        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);
+        for (int force_f32_prec = 0; force_f32_prec < 2; ++force_f32_prec) {
+            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_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++;
+            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);
+            if (force_f32_prec) {
+                ggml_mul_mat_set_prec(out, GGML_PREC_F32);
             }
+
+            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++;
+                }
+            }
+            if (n_mismatch_best < 0 || n_mismatch < n_mismatch_best) {
+                n_mismatch_best = n_mismatch;
+                max_diff_best   = max_diff;
+            }
+
+            ggml_backend_buffer_free(buf);
+            ggml_free(ctx);
         }
-        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) {
+
+        printf("%s: %s GEMM path vs F16 GEMM on fp16-rounded weights (best of default/F32 prec): %d mismatches, max |diff| = %g\n",
+               n_mismatch_best == 0 ? "OK" : "FAILED", ggml_type_name(type), n_mismatch_best, max_diff_best);
+        if (n_mismatch_best != 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 —
@@ -461,21 +504,25 @@ static void build_control_data(ggml_type type, std::vector<uint8_t> & data, std:
 }
 
 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
+    // Only CUDA, HIP (which reports itself as "ROCm") and Vulkan implement
+    // DT3. Any other GPU backend is skipped rather than failed: SYCL answers
+    // supports_op == false for DT3, which is the correct answer for it 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) {
+        // IGPU is a distinct device type from GPU: an integrated Vulkan device
+        // with unified memory reports as IGPU, and accepting only GPU silently
+        // skipped the very hardware this backend is for.
+        const auto dt = ggml_backend_dev_type(dev);
+        if (dt != GGML_BACKEND_DEVICE_TYPE_GPU && dt != GGML_BACKEND_DEVICE_TYPE_IGPU) {
             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);
+        if (strncmp(name, "CUDA", 4) != 0 && strncmp(name, "ROCm", 4) != 0 && strncmp(name, "Vulkan", 6) != 0) {
+            printf("skipping GPU backend %s: DT3 is only implemented for CUDA/HIP/Vulkan\n", name);
             continue;
         }
         backend = ggml_backend_dev_init(dev, nullptr);
@@ -483,7 +530,7 @@ int main(void) {
         break;
     }
     if (backend == nullptr) {
-        printf("no CUDA/HIP backend available, skipping\n");
+        printf("no CUDA/HIP/Vulkan backend available, skipping\n");
         return 0;
     }
 
@@ -497,7 +544,7 @@ int main(void) {
 
     // 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);
+    std::vector<float> y((size_t)NCOLS*100);
     for (size_t i = 0; i < y.size(); ++i) {
         y[i] = i % 32 == 0 ? 127.0f : (float)((int)(rng_next() % 255) - 127);
     }