cuda : add MMQ kernel for DT3

Two decoded ternary planes per SRAM tile row (the planes cannot be fused
into one int8 because d1 != d2), each with its own per-chunk scales, both
multiplied against the same q8_1 y tile: sum = dB*(sumi1*dA1 + sumi2*dA2).
The load decodes each packed byte once with the same base-3 digit
iteration as the MMVQ vec_dot and turns digit bytes {0,1,2} into trit
bytes {-1,0,+1} without cross-byte borrows.

The MMA tile at I=128 takes 592 B/row: 75776 B of x tile plus the y tile,
94720 B at J=128 — fits the 99 KiB opt-in limit of Ampere-class devices
but not e.g. Turing's 64 KiB, so the runtime gate requires the MMA data
layout and enough shared memory for the narrowest tile and declines
otherwise (AMD keeps declining: no config entries select DT3).

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

@@ -13,6 +13,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 +264,18 @@ 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;
+        }
+        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);