summaryrefslogtreecommitdiff
path: root/ggml-cuda.cu
diff options
context:
space:
mode:
Diffstat (limited to 'ggml-cuda.cu')
-rw-r--r--ggml-cuda.cu491
1 files changed, 491 insertions, 0 deletions
diff --git a/ggml-cuda.cu b/ggml-cuda.cu
index 98170a3a..5385e012 100644
--- a/ggml-cuda.cu
+++ b/ggml-cuda.cu
@@ -3,6 +3,7 @@
#include <stdint.h>
#include <stdio.h>
#include <atomic>
+#include <assert.h>
#include <cuda_runtime.h>
#include <cublas_v2.h>
@@ -35,6 +36,7 @@ static_assert(sizeof(half) == sizeof(ggml_fp16_t), "wrong fp16 size");
typedef void (*dequantize_kernel_t)(const void * vx, const int ib, const int iqs, float & v0, float & v1);
typedef void (*to_fp32_cuda_t)(const void * x, float * y, int k, cudaStream_t stream);
typedef void (*dequantize_mul_mat_vec_cuda_t)(const void * vx, const float * y, float * dst, const int ncols, const int nrows, cudaStream_t stream);
+typedef void (*dot_kernel_k_t)(const void * vx, const int ib, const int iqs, const float * y, float & v);
// QK = number of values after dequantization
// QR = QK / number of values before dequantization
@@ -83,6 +85,51 @@ typedef struct {
} block_q8_0;
static_assert(sizeof(block_q8_0) == sizeof(ggml_fp16_t) + QK8_0, "wrong q8_0 block size/padding");
+//================================= k-quants
+
+#define QK_K 256
+
+typedef struct {
+ uint8_t scales[QK_K/16]; // scales and mins, quantized with 4 bits
+ uint8_t qs[QK_K/4]; // quants
+ half d; // super-block scale for quantized scales
+ half dmin; // super-block scale for quantized mins
+} block_q2_k;
+static_assert(sizeof(block_q2_k) == 2*sizeof(ggml_fp16_t) + QK_K/16 + QK_K/4, "wrong q2_k block size/padding");
+
+typedef struct {
+ uint8_t hmask[QK_K/8];
+ uint8_t qs[QK_K/4]; // nibbles / quants
+ uint8_t scales[3*QK_K/64];
+ half d;
+} block_q3_k;
+static_assert(sizeof(block_q3_k) == sizeof(ggml_fp16_t) + QK_K / 4 + 11 * QK_K / 64, "wrong q3_k block size/padding");
+
+typedef struct {
+ half d; // super-block scale for quantized scales
+ half dmin; // super-block scale for quantized mins
+ uint8_t scales[3*QK_K/64]; // scales, quantized with 6 bits
+ uint8_t qs[QK_K/2]; // 4--bit quants
+} block_q4_k;
+static_assert(sizeof(block_q4_k) == 2*sizeof(ggml_fp16_t) + 3*QK_K/64 + QK_K/2, "wrong q4_k block size/padding");
+
+typedef struct {
+ half d; // super-block scale for quantized scales
+ half dmin; // super-block scale for quantized mins
+ uint8_t scales[3*QK_K/64]; // scales, quantized with 6 bits
+ uint8_t qh[QK_K/8]; // quants, high bit
+ uint8_t qs[QK_K/2]; // quants, low 4 bits
+} block_q5_k;
+static_assert(sizeof(block_q5_k) == 2*sizeof(ggml_fp16_t) + 3*QK_K/64 + QK_K/2 + QK_K/8, "wrong q5_k block size/padding");
+
+typedef struct {
+ uint8_t ql[QK_K/2]; // quants, lower 4 bits
+ uint8_t qh[QK_K/4]; // quants, upper 2 bits
+ int8_t scales[QK_K/16]; // scales
+ half d; // delta
+} block_q6_k;
+static_assert(sizeof(block_q6_k) == sizeof(ggml_fp16_t) + 13*QK_K/16, "wrong q6_k block size/padding");
+
#define WARP_SIZE 32
#define CUDA_MUL_BLOCK_SIZE 256
@@ -184,6 +231,337 @@ static __device__ void dequantize_q8_0(const void * vx, const int ib, const int
v1 = vi1*d;
}
+//================================== k-quants
+
+static __global__ void dequantize_block_q2_k(const void * vx, float * yy) {
+
+ const int i = blockIdx.x;
+ const int tid = threadIdx.x;
+ const int n = tid/32;
+ const int l = tid - 32*n;
+ const int is = 8*n + l/16;
+
+ const block_q2_k * x = (const block_q2_k *) vx;
+
+ const uint8_t q = x[i].qs[32*n + l];
+ float * y = yy + i*QK_K + 128*n;
+
+ float dall = x[i].d;
+ float dmin = x[i].dmin;
+ y[l+ 0] = dall * (x[i].scales[is+0] & 0xF) * ((q >> 0) & 3) - dmin * (x[i].scales[is+0] >> 4);
+ y[l+32] = dall * (x[i].scales[is+2] & 0xF) * ((q >> 2) & 3) - dmin * (x[i].scales[is+2] >> 4);
+ y[l+64] = dall * (x[i].scales[is+4] & 0xF) * ((q >> 4) & 3) - dmin * (x[i].scales[is+4] >> 4);
+ y[l+96] = dall * (x[i].scales[is+6] & 0xF) * ((q >> 6) & 3) - dmin * (x[i].scales[is+6] >> 4);
+
+}
+
+static __device__ void vec_dot_q2_k(const void * vx, const int ib, const int iqs, const float * yy, float & result) {
+
+ const block_q2_k * x = (const block_q2_k *) vx;
+
+ // if n is 0, we want to do the lower 128, else the upper 128,
+ // covering y[l+0], y[l+32], y[l+64], y[l+96] and
+ // y[l+16], y[l+48], y[l+80], y[l+112]
+ int n = iqs/128; // 0 or 1
+ int r = iqs - 128*n; // 0...120 in steps of 8
+ int l = r/8; // 0...15 in steps of 1
+
+ const float * y = yy + 128*n + l;
+ const uint8_t * q = x[ib].qs + 32*n + l;
+ const uint8_t * s = x[ib].scales + 8*n;
+
+ const float dall = x[ib].d;
+ const float dmin = x[ib].dmin;
+
+ float sum = y[ 0] * (dall * ((s[0] & 0xF) * ((q[ 0] >> 0) & 3)) - dmin * (s[0] >> 4))
+ + y[ 32] * (dall * ((s[2] & 0xF) * ((q[ 0] >> 2) & 3)) - dmin * (s[2] >> 4))
+ + y[ 64] * (dall * ((s[4] & 0xF) * ((q[ 0] >> 4) & 3)) - dmin * (s[4] >> 4))
+ + y[ 96] * (dall * ((s[6] & 0xF) * ((q[ 0] >> 6) & 3)) - dmin * (s[6] >> 4))
+ + y[ 16] * (dall * ((s[1] & 0xF) * ((q[16] >> 0) & 3)) - dmin * (s[1] >> 4))
+ + y[ 48] * (dall * ((s[3] & 0xF) * ((q[16] >> 2) & 3)) - dmin * (s[3] >> 4))
+ + y[ 80] * (dall * ((s[5] & 0xF) * ((q[16] >> 4) & 3)) - dmin * (s[5] >> 4))
+ + y[112] * (dall * ((s[7] & 0xF) * ((q[16] >> 6) & 3)) - dmin * (s[7] >> 4));
+
+ result = sum;
+
+}
+
+static __global__ void dequantize_block_q3_k(const void * vx, float * yy) {
+
+ int r = threadIdx.x/4;
+ int i = blockIdx.x;
+ int tid = r/2;
+ int is0 = r%2;
+ int l0 = 16*is0 + 4*(threadIdx.x%4);
+ int n = tid / 4;
+ int j = tid - 4*n;
+
+ const block_q3_k * x = (const block_q3_k *) vx;
+
+ uint8_t m = 1 << (4*n + j);
+ int is = 8*n + 2*j + is0;
+ int shift = 2*j;
+
+ int8_t us = is < 4 ? (x[i].scales[is-0] & 0xF) | (((x[i].scales[is+8] >> 0) & 3) << 4) :
+ is < 8 ? (x[i].scales[is-0] & 0xF) | (((x[i].scales[is+4] >> 2) & 3) << 4) :
+ is < 12 ? (x[i].scales[is-8] >> 4) | (((x[i].scales[is+0] >> 4) & 3) << 4) :
+ (x[i].scales[is-8] >> 4) | (((x[i].scales[is-4] >> 6) & 3) << 4);
+ float d_all = x[i].d;
+ float dl = d_all * (us - 32);
+
+ float * y = yy + i*QK_K + 128*n + 32*j;
+ const uint8_t * q = x[i].qs + 32*n;
+ const uint8_t * hm = x[i].hmask;
+
+ for (int l = l0; l < l0+4; ++l) y[l] = dl * ((int8_t)((q[l] >> shift) & 3) - ((hm[l] & m) ? 0 : 4));
+
+}
+
+static __device__ void vec_dot_q3_k(const void * vx, const int ib, const int iqs, const float * yy, float & result) {
+
+ const block_q3_k * x = (const block_q3_k *) vx;
+
+ const uint32_t kmask1 = 0x03030303;
+ const uint32_t kmask2 = 0x0f0f0f0f;
+
+ uint32_t aux[3];
+ uint32_t utmp[4];
+
+ // if n is 0, we want to do the lower 128, else the upper 128,
+ // covering y[l+0], y[l+32], y[l+64], y[l+96] and
+ // y[l+16], y[l+48], y[l+80], y[l+112]
+ int n = iqs/128; // 0 or 1
+ int r = iqs - 128*n; // 0...120 in steps of 8
+ int l = r/8; // 0...15 in steps of 1
+
+ const float * y = yy + 128*n + l;
+ const uint8_t * q = x[ib].qs + 32*n + l;
+ const uint8_t * hm = x[ib].hmask + l;
+ const int8_t * s = (const int8_t *)utmp + 8*n;
+
+ memcpy(aux, x[ib].scales, 12);
+ utmp[3] = ((aux[1] >> 4) & kmask2) | (((aux[2] >> 6) & kmask1) << 4);
+ utmp[2] = ((aux[0] >> 4) & kmask2) | (((aux[2] >> 4) & kmask1) << 4);
+ utmp[1] = (aux[1] & kmask2) | (((aux[2] >> 2) & kmask1) << 4);
+ utmp[0] = (aux[0] & kmask2) | (((aux[2] >> 0) & kmask1) << 4);
+
+ const float dall = x[ib].d;
+
+ const uint8_t m = 1 << (4*n);
+
+ float sum = y[ 0] * (s[0] - 32) * (((q[ 0] >> 0) & 3) - (hm[ 0] & (m << 0) ? 0 : 4))
+ + y[ 32] * (s[2] - 32) * (((q[ 0] >> 2) & 3) - (hm[ 0] & (m << 1) ? 0 : 4))
+ + y[ 64] * (s[4] - 32) * (((q[ 0] >> 4) & 3) - (hm[ 0] & (m << 2) ? 0 : 4))
+ + y[ 96] * (s[6] - 32) * (((q[ 0] >> 6) & 3) - (hm[ 0] & (m << 3) ? 0 : 4))
+ + y[ 16] * (s[1] - 32) * (((q[16] >> 0) & 3) - (hm[16] & (m << 0) ? 0 : 4))
+ + y[ 48] * (s[3] - 32) * (((q[16] >> 2) & 3) - (hm[16] & (m << 1) ? 0 : 4))
+ + y[ 80] * (s[5] - 32) * (((q[16] >> 4) & 3) - (hm[16] & (m << 2) ? 0 : 4))
+ + y[112] * (s[7] - 32) * (((q[16] >> 6) & 3) - (hm[16] & (m << 3) ? 0 : 4));
+
+ result = sum * dall;
+
+}
+
+static inline __device__ void get_scale_min_k4(int j, const uint8_t * q, uint8_t & d, uint8_t & m) {
+ if (j < 4) {
+ d = q[j] & 63; m = q[j + 4] & 63;
+ } else {
+ d = (q[j+4] & 0xF) | ((q[j-4] >> 6) << 4);
+ m = (q[j+4] >> 4) | ((q[j-0] >> 6) << 4);
+ }
+}
+
+static __global__ void dequantize_block_q4_k(const void * vx, float * yy) {
+ const block_q4_k * x = (const block_q4_k *) vx;
+
+ const int i = blockIdx.x;
+
+ //// assume 64 threads - this is very slightly better than the one below
+ //const int tid = threadIdx.x;
+ //const int il = tid/16;
+ //const int ir = tid%16;
+ //const int is = 2*il;
+ //const int n = 2;
+
+ // assume 32 threads
+ const int tid = threadIdx.x;
+ const int il = tid/8;
+ const int ir = tid%8;
+ const int is = 2*il;
+ const int n = 4;
+
+ float * y = yy + i*QK_K + 64*il + n*ir;
+
+ const float dall = x[i].d;
+ const float dmin = x[i].dmin;
+
+ const uint8_t * q = x[i].qs + 32*il + n*ir;
+
+ uint8_t sc, m;
+ get_scale_min_k4(is + 0, x[i].scales, sc, m);
+ const float d1 = dall * sc; const float m1 = dmin * m;
+ get_scale_min_k4(is + 1, x[i].scales, sc, m);
+ const float d2 = dall * sc; const float m2 = dmin * m;
+ for (int l = 0; l < n; ++l) {
+ y[l + 0] = d1 * (q[l] & 0xF) - m1;
+ y[l +32] = d2 * (q[l] >> 4) - m2;
+ }
+}
+
+static __device__ void vec_dot_q4_k(const void * vx, const int ib, const int iqs, const float * yy, float & result) {
+
+ const block_q4_k * x = (const block_q4_k *) vx;
+
+ // iqs is in 0...248 in steps of 8 =>
+ const int j = iqs / 64; // j is in 0...3
+ const int ir = (iqs - 64*j)/2; // ir is in 0...28 in steps of 4
+ const int is = 2*j; // is is in 0...6 in steps of 2
+
+ const float * y = yy + 64*j + ir;
+ const uint8_t * q = x[ib].qs + 32*j + ir;
+
+ const float dall = x[ib].d;
+ const float dmin = x[ib].dmin;
+
+ uint8_t sc, m;
+ get_scale_min_k4(is + 0, x[ib].scales, sc, m);
+ const float d1 = dall * sc;
+ const float m1 = dmin * m;
+ get_scale_min_k4(is + 1, x[ib].scales, sc, m);
+ const float d2 = dall * sc;
+ const float m2 = dmin * m;
+
+ float sum = 0;
+ for (int k = 0; k < 4; ++k) {
+ sum += y[k + 0] * (d1 * (q[k] & 0xF) - m1);
+ sum += y[k + 32] * (d2 * (q[k] >> 4) - m2);
+ }
+ result = sum;
+
+}
+
+static __global__ void dequantize_block_q5_k(const void * vx, float * yy) {
+ const block_q5_k * x = (const block_q5_k *) vx;
+
+ const int i = blockIdx.x;
+
+ // assume 64 threads - this is very slightly better than the one below
+ const int tid = threadIdx.x;
+ const int il = tid/16; // il is in 0...3
+ const int ir = tid%16; // ir is in 0...15
+ const int is = 2*il; // is is in 0...6
+
+ float * y = yy + i*QK_K + 64*il + 2*ir;
+
+ const float dall = x[i].d;
+ const float dmin = x[i].dmin;
+
+ const uint8_t * ql = x[i].qs + 32*il + 2*ir;
+ const uint8_t * qh = x[i].qh + 2*ir;
+
+ uint8_t sc, m;
+ get_scale_min_k4(is + 0, x[i].scales, sc, m);
+ const float d1 = dall * sc; const float m1 = dmin * m;
+ get_scale_min_k4(is + 1, x[i].scales, sc, m);
+ const float d2 = dall * sc; const float m2 = dmin * m;
+
+ uint8_t hm = 1 << (2*il);
+ y[ 0] = d1 * ((ql[ 0] & 0xF) + (qh[ 0] & hm ? 16 : 0)) - m1;
+ y[ 1] = d1 * ((ql[ 1] & 0xF) + (qh[ 1] & hm ? 16 : 0)) - m1;
+ hm <<= 1;
+ y[32] = d2 * ((ql[ 0] >> 4) + (qh[ 0] & hm ? 16 : 0)) - m2;
+ y[33] = d2 * ((ql[ 1] >> 4) + (qh[ 1] & hm ? 16 : 0)) - m2;
+}
+
+static __device__ void vec_dot_q5_k(const void * vx, const int ib, const int iqs, const float * yy, float & result) {
+
+ const block_q5_k * x = (const block_q5_k *) vx;
+
+ // iqs is in 0...248 in steps of 8 =>
+ const int j = iqs / 64; // j is in 0...3
+ const int ir = (iqs - 64*j)/2; // ir is in 0...28 in steps of 4
+ const int is = 2*j; // is is in 0...6 in steps of 2
+
+ const float * y = yy + 64*j + ir;
+ const uint8_t * ql = x[ib].qs + 32*j + ir;
+ const uint8_t * qh = x[ib].qh + ir;
+
+ const float dall = x[ib].d;
+ const float dmin = x[ib].dmin;
+
+ uint8_t sc, m;
+ get_scale_min_k4(is + 0, x[ib].scales, sc, m);
+ const float d1 = dall * sc;
+ const float m1 = dmin * m;
+ get_scale_min_k4(is + 1, x[ib].scales, sc, m);
+ const float d2 = dall * sc;
+ const float m2 = dmin * m;
+
+ uint8_t hm = 1 << is;
+ float sum = 0;
+ for (int k = 0; k < 4; ++k) {
+ sum += y[k + 0] * (d1 * ((ql[k] & 0xF) + (qh[k] & hm ? 16 : 0)) - m1);
+ }
+ hm <<= 1;
+ for (int k = 0; k < 4; ++k) {
+ sum += y[k + 32] * (d2 * ((ql[k] >> 4) + (qh[k] & hm ? 16 : 0)) - m2);
+ }
+ result = sum;
+
+}
+
+static __global__ void dequantize_block_q6_k(const void * vx, float * yy) {
+ const block_q6_k * x = (const block_q6_k *) vx;
+
+ const int i = blockIdx.x;
+
+ // assume 64 threads - this is very slightly better than the one below
+ const int tid = threadIdx.x;
+ const int ip = tid/32; // ip is 0 or 1
+ const int il = tid - 32*ip; // 0...32
+ const int is = 8*ip + il/16;
+
+ float * y = yy + i*QK_K + 128*ip + il;
+
+ const float d = x[i].d;
+
+ const uint8_t * ql = x[i].ql + 64*ip + il;
+ const uint8_t qh = x[i].qh[32*ip + il];
+ const int8_t * sc = x[i].scales + is;
+
+ y[ 0] = d * sc[0] * ((int8_t)((ql[ 0] & 0xF) | (((qh >> 0) & 3) << 4)) - 32);
+ y[32] = d * sc[2] * ((int8_t)((ql[32] & 0xF) | (((qh >> 2) & 3) << 4)) - 32);
+ y[64] = d * sc[4] * ((int8_t)((ql[ 0] >> 4) | (((qh >> 4) & 3) << 4)) - 32);
+ y[96] = d * sc[6] * ((int8_t)((ql[32] >> 4) | (((qh >> 6) & 3) << 4)) - 32);
+}
+
+static __device__ void vec_dot_q6_k(const void * vx, const int ib, const int iqs, const float * yy, float & result) {
+
+ const block_q6_k * x = (const block_q6_k *) vx;
+
+ const int ip = iqs / 128; // 0 or 1
+ const int il = (iqs - 128*ip)/8; // 0...15
+ const int is = 8*ip;
+
+ const float * y = yy + 128*ip + il;
+
+ const float d = x[ib].d;
+
+ const uint8_t * ql = x[ib].ql + 64*ip + il;
+ const uint8_t * qh = x[ib].qh + 32*ip + il;
+ const int8_t * sc = x[ib].scales + is;
+
+ result = y[ 0] * d * sc[0] * ((int8_t)((ql[ 0] & 0xF) | (((qh[ 0] >> 0) & 3) << 4)) - 32)
+ + y[ 32] * d * sc[2] * ((int8_t)((ql[32] & 0xF) | (((qh[ 0] >> 2) & 3) << 4)) - 32)
+ + y[ 64] * d * sc[4] * ((int8_t)((ql[ 0] >> 4) | (((qh[ 0] >> 4) & 3) << 4)) - 32)
+ + y[ 96] * d * sc[6] * ((int8_t)((ql[32] >> 4) | (((qh[ 0] >> 6) & 3) << 4)) - 32)
+ + y[ 16] * d * sc[1] * ((int8_t)((ql[16] & 0xF) | (((qh[16] >> 0) & 3) << 4)) - 32)
+ + y[ 48] * d * sc[3] * ((int8_t)((ql[48] & 0xF) | (((qh[16] >> 2) & 3) << 4)) - 32)
+ + y[ 80] * d * sc[5] * ((int8_t)((ql[16] >> 4) | (((qh[16] >> 4) & 3) << 4)) - 32)
+ + y[112] * d * sc[7] * ((int8_t)((ql[48] >> 4) | (((qh[16] >> 6) & 3) << 4)) - 32);
+
+}
+
static __device__ void convert_f16(const void * vx, const int ib, const int iqs, float & v0, float & v1){
const half * x = (const half *) vx;
@@ -258,6 +636,41 @@ static __global__ void dequantize_mul_mat_vec(const void * vx, const float * y,
}
}
+template <int n_thread, dot_kernel_k_t dot_kernel>
+static __global__ void dequantize_mul_mat_vec_k(const void * vx, const float * y, float * dst, const int ncols) {
+ const int row = blockIdx.x*blockDim.y + threadIdx.y;
+ const int tid = threadIdx.x;
+
+ const int iter_stride = QK_K;
+ const int vals_per_iter = iter_stride / n_thread;
+ const int num_blocks_per_row = ncols / QK_K;
+ const int ib0 = row*num_blocks_per_row;
+
+ float tmp = 0; // partial sum for thread in warp
+
+ for (int i = 0; i < ncols; i += iter_stride) {
+ const int col = i + vals_per_iter*tid;
+ const int ib = ib0 + col/QK_K; // x block index
+ const int iqs = col%QK_K; // x quant index
+ const int iybs = col - col%QK_K; // y block start index
+
+ float v;
+ dot_kernel(vx, ib, iqs, y + iybs, v);
+ tmp += v;
+ }
+
+ // sum up partial sums and write back result
+ __syncthreads();
+#pragma unroll
+ for (int mask = 16; mask > 0; mask >>= 1) {
+ tmp += __shfl_xor_sync(0xffffffff, tmp, mask, 32);
+ }
+
+ if (tid == 0) {
+ dst[row] = tmp;
+ }
+}
+
static void mul_f32_cuda(const float * x, const float * y, float * dst, const int kx, const int ky, cudaStream_t stream) {
const int num_blocks = (kx + CUDA_MUL_BLOCK_SIZE - 1) / CUDA_MUL_BLOCK_SIZE;
mul_f32<<<num_blocks, CUDA_MUL_BLOCK_SIZE, 0, stream>>>(x, y, dst, kx, ky);
@@ -288,6 +701,31 @@ static void dequantize_row_q8_0_cuda(const void * vx, float * y, const int k, cu
dequantize_block<QK8_0, QR8_0, dequantize_q8_0><<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>>(vx, y, k);
}
+static void dequantize_row_q2_k_cuda(const void * vx, float * y, const int k, cudaStream_t stream) {
+ const int nb = k / QK_K;
+ dequantize_block_q2_k<<<nb, 64, 0, stream>>>(vx, y);
+}
+
+static void dequantize_row_q3_k_cuda(const void * vx, float * y, const int k, cudaStream_t stream) {
+ const int nb = k / QK_K;
+ dequantize_block_q3_k<<<nb, 64, 0, stream>>>(vx, y);
+}
+
+static void dequantize_row_q4_k_cuda(const void * vx, float * y, const int k, cudaStream_t stream) {
+ const int nb = k / QK_K;
+ dequantize_block_q4_k<<<nb, 32, 0, stream>>>(vx, y);
+}
+
+static void dequantize_row_q5_k_cuda(const void * vx, float * y, const int k, cudaStream_t stream) {
+ const int nb = k / QK_K;
+ dequantize_block_q5_k<<<nb, 64, 0, stream>>>(vx, y);
+}
+
+static void dequantize_row_q6_k_cuda(const void * vx, float * y, const int k, cudaStream_t stream) {
+ const int nb = k / QK_K;
+ dequantize_block_q6_k<<<nb, 64, 0, stream>>>(vx, y);
+}
+
static void dequantize_mul_mat_vec_q4_0_cuda(const void * vx, const float * y, float * dst, const int ncols, const int nrows, cudaStream_t stream) {
GGML_ASSERT(ncols % GGML_CUDA_DMMV_X == 0);
GGML_ASSERT(nrows % GGML_CUDA_DMMV_Y == 0);
@@ -328,6 +766,37 @@ static void dequantize_mul_mat_vec_q8_0_cuda(const void * vx, const float * y, f
<<<nrows/GGML_CUDA_DMMV_Y, block_dims, 0, stream>>>(vx, y, dst, ncols);
}
+static void dequantize_mul_mat_vec_q2_k_cuda(const void * vx, const float * y, float * dst, const int ncols, const int nrows, cudaStream_t stream) {
+ GGML_ASSERT(ncols % QK_K == 0);
+ const int ny = 2;
+ const dim3 block_dims(32, ny, 1);
+ dequantize_mul_mat_vec_k<32, vec_dot_q2_k><<<(nrows + ny - 1)/ny, block_dims, 0, stream>>>(vx, y, dst, ncols);
+}
+
+static void dequantize_mul_mat_vec_q3_k_cuda(const void * vx, const float * y, float * dst, const int ncols, const int nrows, cudaStream_t stream) {
+ GGML_ASSERT(ncols % QK_K == 0);
+ const dim3 block_dims(32, 2, 1);
+ dequantize_mul_mat_vec_k<32, vec_dot_q3_k><<<nrows/2, block_dims, 0, stream>>>(vx, y, dst, ncols);
+}
+
+static void dequantize_mul_mat_vec_q4_k_cuda(const void * vx, const float * y, float * dst, const int ncols, const int nrows, cudaStream_t stream) {
+ GGML_ASSERT(ncols % QK_K == 0);
+ const dim3 block_dims(32, 2, 1);
+ dequantize_mul_mat_vec_k<32, vec_dot_q4_k><<<nrows/2, block_dims, 0, stream>>>(vx, y, dst, ncols);
+}
+
+static void dequantize_mul_mat_vec_q5_k_cuda(const void * vx, const float * y, float * dst, const int ncols, const int nrows, cudaStream_t stream) {
+ GGML_ASSERT(ncols % QK_K == 0);
+ const dim3 block_dims(32, 2, 1);
+ dequantize_mul_mat_vec_k<32, vec_dot_q5_k><<<nrows/2, block_dims, 0, stream>>>(vx, y, dst, ncols);
+}
+
+static void dequantize_mul_mat_vec_q6_k_cuda(const void * vx, const float * y, float * dst, const int ncols, const int nrows, cudaStream_t stream) {
+ GGML_ASSERT(ncols % QK_K == 0);
+ const dim3 block_dims(32, 2, 1);
+ dequantize_mul_mat_vec_k<32, vec_dot_q6_k><<<nrows/2, block_dims, 0, stream>>>(vx, y, dst, ncols);
+}
+
static void convert_fp16_to_fp32_cuda(const void * vx, float * y, const int k, cudaStream_t stream) {
const int num_blocks = (k + CUDA_DEQUANTIZE_BLOCK_SIZE - 1) / CUDA_DEQUANTIZE_BLOCK_SIZE;
dequantize_block<32, 1, convert_f16><<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>>(vx, y, k);
@@ -353,6 +822,16 @@ static to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
return dequantize_row_q5_1_cuda;
case GGML_TYPE_Q8_0:
return dequantize_row_q8_0_cuda;
+ case GGML_TYPE_Q2_K:
+ return dequantize_row_q2_k_cuda;
+ case GGML_TYPE_Q3_K:
+ return dequantize_row_q3_k_cuda;
+ case GGML_TYPE_Q4_K:
+ return dequantize_row_q4_k_cuda;
+ case GGML_TYPE_Q5_K:
+ return dequantize_row_q5_k_cuda;
+ case GGML_TYPE_Q6_K:
+ return dequantize_row_q6_k_cuda;
case GGML_TYPE_F16:
return convert_fp16_to_fp32_cuda;
default:
@@ -372,6 +851,16 @@ static dequantize_mul_mat_vec_cuda_t ggml_get_dequantize_mul_mat_vec_cuda(ggml_t
return dequantize_mul_mat_vec_q5_1_cuda;
case GGML_TYPE_Q8_0:
return dequantize_mul_mat_vec_q8_0_cuda;
+ case GGML_TYPE_Q2_K:
+ return dequantize_mul_mat_vec_q2_k_cuda;
+ case GGML_TYPE_Q3_K:
+ return dequantize_mul_mat_vec_q3_k_cuda;
+ case GGML_TYPE_Q4_K:
+ return dequantize_mul_mat_vec_q4_k_cuda;
+ case GGML_TYPE_Q5_K:
+ return dequantize_mul_mat_vec_q5_k_cuda;
+ case GGML_TYPE_Q6_K:
+ return dequantize_mul_mat_vec_q6_k_cuda;
case GGML_TYPE_F16:
return convert_mul_mat_vec_f16_cuda;
default:
@@ -790,12 +1279,14 @@ static void ggml_cuda_mul_mat_q_f32(const ggml_tensor * src0, const ggml_tensor
CUDA_CHECK(cudaStreamWaitEvent(cudaStream, cudaEvent, 0));
// compute
+ //printf("Calling dmmv\n");
dmmv(c_Q, c_Y, c_D, ne00, ne01, cudaStream);
CUDA_CHECK(cudaGetLastError());
} else { // general dequantization kernel + cuBLAS matrix matrix multiplication
float * c_X = d_X + i * x_ne;
+//typedef void (*to_fp32_cuda_t)(const void * x, float * y, int k, cudaStream_t stream);
// convert src0 to fp32 on device
to_fp32_cuda(c_Q, c_X, x_ne, cudaStream2);
CUDA_CHECK(cudaGetLastError());
generated by cgit v1.2.3 (git 2.25.1) at 2025-08-16 01:12:43 +0000