diff options
author | Stephan Walter <stephan@walter.name> | 2023-04-28 23:10:43 +0000 |
---|---|---|
committer | GitHub <noreply@github.com> | 2023-04-28 23:10:43 +0000 |
commit | 36d19a603b221d1bd7897fcb10e823e2103b052d (patch) | |
tree | 14e8caf8b61f8fa0fa72f5d7cc137b1ff9330b2f /ggml-cuda.cu | |
parent | 7f15c5c477d9933689a9d1c40794483e350c2f19 (diff) |
Remove Q4_3 which is no better than Q5 (#1218)
Diffstat (limited to 'ggml-cuda.cu')
-rw-r--r-- | ggml-cuda.cu | 37 |
1 files changed, 0 insertions, 37 deletions
diff --git a/ggml-cuda.cu b/ggml-cuda.cu index b1bd29b..d619f5d 100644 --- a/ggml-cuda.cu +++ b/ggml-cuda.cu @@ -29,14 +29,6 @@ typedef struct { } block_q4_2; static_assert(sizeof(block_q4_2) == sizeof(ggml_fp16_t) + QK4_2 / 2, "wrong q4_2 block size/padding"); -#define QK4_3 16 -typedef struct { - __half d; // delta - __half m; // min - uint8_t qs[QK4_3 / 2]; // nibbles / quants -} block_q4_3; -static_assert(sizeof(block_q4_3) == 2 * sizeof(ggml_fp16_t) + QK4_3 / 2, "wrong q4_3 block size/padding"); - #define QK5_0 32 typedef struct { __half d; // delta @@ -131,30 +123,6 @@ static __global__ void dequantize_block_q4_2(const void * vx, float * y) { } } -static __global__ void dequantize_block_q4_3(const void * vx, float * y) { - const block_q4_3 * x = (const block_q4_3 *) vx; - - const int i = blockIdx.x; - - const float d = x[i].d; - const float m = x[i].m; - - const uint8_t * pp = x[i].qs; - - for (int l = 0; l < QK4_3; l += 2) { - const uint8_t vi = pp[l/2]; - - const int8_t vi0 = vi & 0xf; - const int8_t vi1 = vi >> 4; - - const float v0 = vi0*d + m; - const float v1 = vi1*d + m; - - y[i*QK4_3 + l + 0] = v0; - y[i*QK4_3 + l + 1] = v1; - } -} - static __global__ void dequantize_block_q5_0(const void * vx, float * y) { const block_q5_0 * x = (const block_q5_0 *) vx; @@ -244,11 +212,6 @@ void dequantize_row_q4_2_cuda(const void * vx, float * y, int k, cudaStream_t st dequantize_block_q4_2<<<nb, 1, 0, stream>>>(vx, y); } -void dequantize_row_q4_3_cuda(const void * vx, float * y, int k, cudaStream_t stream) { - const int nb = k / QK4_3; - dequantize_block_q4_3<<<nb, 1, 0, stream>>>(vx, y); -} - void dequantize_row_q5_0_cuda(const void * vx, float * y, int k, cudaStream_t stream) { const int nb = k / QK5_0; dequantize_block_q5_0<<<nb, 1, 0, stream>>>(vx, y); |