q5_0

Author: JohannesGaessler

ggml-cuda.cu

@@ -1220,7 +1220,6 @@ static __device__ __forceinline__ float vec_dot_q4_0_q8_1(const void * vbq, cons
     sumi     = __dp4a(vi1, ui1, sumi);
 
     return sumi*d;
-
 }
 
 static __device__ __forceinline__ float vec_dot_q4_1_q8_1(const void * vbq, const block_q8_1 * bq8_1, const int iqs) {
@@ -1241,7 +1240,37 @@ static __device__ __forceinline__ float vec_dot_q4_1_q8_1(const void * vbq, cons
     sumi     = __dp4a(vi1, ui1, sumi);
 
     return sumi*d + m*s / QI4_1;
+}
+
+static __device__ __forceinline__ float vec_dot_q5_0_q8_1(const void * vbq, const block_q8_1 * bq8_1, const int iqs) {
+    const block_q5_0 * bq5_0 = (const block_q5_0 *) vbq;
+
+    int qs;
+    memcpy(&qs, &bq5_0->qs[sizeof(int) * (iqs + 0)], sizeof(int));
+    const int qh0 = bq5_0->qh[iqs/2 + 0] >> 4*(iqs%2);
+    const int qh1 = bq5_0->qh[iqs/2 + 2] >> 4*(iqs%2);
+    const int ui0 = *((int *) &bq8_1->qs[sizeof(int) * (iqs + 0)]);
+    const int ui1 = *((int *) &bq8_1->qs[sizeof(int) * (iqs + QI4_0)]);
+
+    const float d = bq5_0->d * bq8_1->d;
+
+    int vi0 = (qs  >>  0) & 0x0F0F0F0F;
+    vi0    |= (qh0 <<  4) & 0x00000010;
+    vi0    |= (qh0 << 11) & 0x00001000;
+    vi0    |= (qh0 << 18) & 0x00100000;
+    vi0    |= (qh0 << 25) & 0x10000000;
+    vi0     = __vsub4(vi0,  0x10101010);
+    int sumi = __dp4a(vi0, ui0, 0);
+
+    int vi1 = (qs  >>  4) & 0x0F0F0F0F;
+    vi1    |= (qh1 <<  4) & 0x00000010;
+    vi1    |= (qh1 << 11) & 0x00001000;
+    vi1    |= (qh1 << 18) & 0x00100000;
+    vi1    |= (qh1 << 25) & 0x10000000;
+    vi1     = __vsub4(vi1,  0x10101010);
+    sumi = __dp4a(vi1, ui1, sumi);
 
+    return sumi*d;
 }
 
 template <int qk, int qr, dequantize_kernel_t dequantize_kernel>
@@ -1796,6 +1825,15 @@ static void mul_mat_vec_q4_1_q8_1_cuda(const void * vx, const void * vy, float *
         <<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
 }
 
+static void mul_mat_vec_q5_0_q8_1_cuda(const void * vx, const void * vy, float * dst, const int ncols, const int nrows, cudaStream_t stream) {
+    GGML_ASSERT(ncols % GGML_CUDA_DMMV_X == 0);
+    const int block_num_y = (nrows + GGML_CUDA_DMMV_Y - 1) / GGML_CUDA_DMMV_Y;
+    const dim3 block_nums(1, block_num_y, 1);
+    const dim3 block_dims(WARP_SIZE, GGML_CUDA_DMMV_Y, 1);
+    mul_mat_vec_q<QK5_0, block_q5_0, vec_dot_q5_0_q8_1>
+        <<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
+}
+
 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<1, 1, convert_f16><<<num_blocks, CUDA_DEQUANTIZE_BLOCK_SIZE, 0, stream>>>(vx, y, k);
@@ -2319,6 +2357,9 @@ inline void ggml_cuda_op_mul_mat_vec_q(
         case GGML_TYPE_Q4_1:
             mul_mat_vec_q4_1_q8_1_cuda(src0_ddq_i, src1_q8_0, dst_ddf_i, ne00, nrows, cudaStream_main);
             break;
+        case GGML_TYPE_Q5_0:
+            mul_mat_vec_q5_0_q8_1_cuda(src0_ddq_i, src1_q8_0, dst_ddf_i, ne00, nrows, cudaStream_main);
+            break;
         default:
             GGML_ASSERT(false);
             break;
@@ -2875,7 +2916,7 @@ void ggml_cuda_mul_mat(const ggml_tensor * src0, const ggml_tensor * src1, ggml_
     } else if (ggml_is_quantized(src0->type) || src0->type == GGML_TYPE_F16) {
         if (src1->ne[1] == 1 && src0->ne[0] % GGML_CUDA_DMMV_X == 0 && src0->ne[1] % GGML_CUDA_DMMV_Y == 0) {
             bool use_mul_mat_vec_q = false;
-            use_mul_mat_vec_q = src0->type == GGML_TYPE_Q4_0 || src0->type == GGML_TYPE_Q4_1;
+            use_mul_mat_vec_q = src0->type == GGML_TYPE_Q4_0 || src0->type == GGML_TYPE_Q4_1 || src0->type == GGML_TYPE_Q5_0;
             if (use_mul_mat_vec_q) {
                 ggml_cuda_op(src0, src1, dst, ggml_cuda_op_mul_mat_vec_q, false, false);
             } else {