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);