More GPU threads for dequantization

Author: JohannesGaessler

ggml-cuda.cu

@@ -83,7 +83,7 @@ static_assert(sizeof(block_q8_0) == sizeof(float) + QK8_0, "wrong q8_0 block siz
 static __global__ void dequantize_block_q4_0(const void * vx, float * y) {
     const block_q4_0 * x = (const block_q4_0 *) vx;
 
-    const int i = blockIdx.x;
+    const int i = blockIdx.x*blockDim.x + threadIdx.x;
 
     const float d = x[i].d;
 
@@ -182,7 +182,7 @@ static __global__ void dequantize_block_q5_0(const void * vx, float * y) {
 static __global__ void dequantize_block_q5_1(const void * vx, float * y) {
     const block_q5_1 * x = (const block_q5_1 *) vx;
 
-    const int i = blockIdx.x;
+    const int i = blockIdx.x*blockDim.x + threadIdx.x;
 
     const float d = x[i].d;
     const float m = x[i].m;
@@ -227,7 +227,8 @@ static __global__ void dequantize_block_q8_0(const void * vx, float * y) {
 
 static void dequantize_row_q4_0_cuda(const void * vx, float * y, int k, cudaStream_t stream) {
     const int nb = k / QK4_0;
-    dequantize_block_q4_0<<<nb, 1, 0, stream>>>(vx, y);
+    GGML_ASSERT(nb % 256 == 0);
+    dequantize_block_q4_0<<<nb/256, 256, 0, stream>>>(vx, y);
 }
 
 static void dequantize_row_q4_1_cuda(const void * vx, float * y, int k, cudaStream_t stream) {
@@ -247,7 +248,8 @@ static void dequantize_row_q5_0_cuda(const void * vx, float * y, int k, cudaStre
 
 static void dequantize_row_q5_1_cuda(const void * vx, float * y, int k, cudaStream_t stream) {
     const int nb = k / QK5_1;
-    dequantize_block_q5_1<<<nb, 1, 0, stream>>>(vx, y);
+    GGML_ASSERT(nb % 256 == 0);
+    dequantize_block_q5_1<<<nb/256, 256, 0, stream>>>(vx, y);
 }
 
 static void dequantize_row_q8_0_cuda(const void * vx, float * y, int k, cudaStream_t stream) {