From f46e8072d8e6c69e6202b1fa764d03bb65c72e7a Mon Sep 17 00:00:00 2001 From: Millaguie Date: Tue, 11 Aug 2026 08:33:06 +0200 Subject: [PATCH] cuda : add MMQ kernel for DT3 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 | 17 +++ ggml/src/ggml-cuda/mmq-load-tiles.cuh | 111 +++++++++++++++ ggml/src/ggml-cuda/mmq-vec-dot.cuh | 134 ++++++++++++++++++ ggml/src/ggml-cuda/mmq.cu | 15 ++ ggml/src/ggml-cuda/mmq.cuh | 20 +++ .../template-instances/generate_cu_files.py | 1 + .../template-instances/mmq-instance-dt3.cu | 5 + 7 files changed, 303 insertions(+) create mode 100644 ggml/src/ggml-cuda/template-instances/mmq-instance-dt3.cu diff --git a/ggml/src/ggml-cuda/mmq-config-ampere.cuh b/ggml/src/ggml-cuda/mmq-config-ampere.cuh index 9f9fd1973..c2556826d 100644 --- a/ggml/src/ggml-cuda/mmq-config-ampere.cuh +++ b/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); diff --git a/ggml/src/ggml-cuda/mmq-load-tiles.cuh b/ggml/src/ggml-cuda/mmq-load-tiles.cuh index 8ed704c28..833a834ec 100644 --- a/ggml/src/ggml-cuda/mmq-load-tiles.cuh +++ b/ggml/src/ggml-cuda/mmq-load-tiles.cuh @@ -176,6 +176,117 @@ template static __device__ __forceinline_ } } +template 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 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(); diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh index d57343386..995242e57 100644 --- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh +++ b/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 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 + (&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 + (&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 +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 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(); diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 707437ea3..bfcf155bc 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/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(ctx, args, stream); break; + case GGML_TYPE_DT3: + mul_mat_q_case(ctx, args, stream); + break; case GGML_TYPE_Q4_0: mul_mat_q_case(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) { diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 2eb15fdfa..288525781 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/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, ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a, ggml_cuda_mmq_write_back_dp4a); + case GGML_TYPE_DT3: + return ggml_cuda_mmq_util_funcs( + VDR_Q8_0_Q8_1_MMQ, + ggml_cuda_mmq_load_tiles_dt3, + ggml_cuda_mmq_vec_dot_dt3_q8_1_dp4a, + ggml_cuda_mmq_write_back_dp4a); 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, ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma, ggml_cuda_mmq_write_back_mma); + case GGML_TYPE_DT3: + return ggml_cuda_mmq_util_funcs( + -1, + ggml_cuda_mmq_load_tiles_dt3, + ggml_cuda_mmq_vec_dot_dt3_q8_1_mma, + ggml_cuda_mmq_write_back_mma); 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); diff --git a/ggml/src/ggml-cuda/template-instances/generate_cu_files.py b/ggml/src/ggml-cuda/template-instances/generate_cu_files.py index d7cd27167..f5853a231 100755 --- a/ggml/src/ggml-cuda/template-instances/generate_cu_files.py +++ b/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", diff --git a/ggml/src/ggml-cuda/template-instances/mmq-instance-dt3.cu b/ggml/src/ggml-cuda/template-instances/mmq-instance-dt3.cu new file mode 100644 index 000000000..c797eb2c3 --- /dev/null +++ b/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);