A high-throughput and memory-efficient inference and serving engine for LLMs
Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git
Synced Nov 15, 19:45 UTC
[Kernel] Add more dtype support for GGUF kernels (#14043)
Signed-off-by: SzymonOzog <szymon.ozog@aleph-alpha.com> Signed-off-by: SzymonOzog <szymon.ozog@gmail.com>
Szymon Ożóg committed
Mar 10, 2025 at 15:30 UTC
89cdaa83e7d049fba4ee6593047f625fa5c0855a
6 files changed
+319
-267
csrc/quantization/gguf/gguf_kernel.cu
+168
-152
@@ -5,6 +5,7 @@
5
#include <c10/cuda/CUDAGuard.h>
6
7
#include "cuda_compat.h"
8
+#include "dispatch_utils.h"
9
10
#include "ggml-common.h"
11
#include "vecdotq.cuh"
@@ -13,7 +14,8 @@
14
#include "mmq.cuh"
15
16
// Q8 gemv
16
-static __global__ void quantize_q8_1(const half* __restrict__ x,
17
+template <typename scalar_t>
18
+static __global__ void quantize_q8_1(const scalar_t* __restrict__ x,
19
void* __restrict__ vy, const int kx,
20
const int kx_padded) {
21
const int ix = blockDim.x * blockIdx.x + threadIdx.x;
@@ -28,7 +30,7 @@ static __global__ void quantize_q8_1(const half* __restrict__ x,
30
const int ib = i_padded / QK8_1; // block index
31
const int iqs = i_padded % QK8_1; // quant index
32
31
- const float xi = ix < kx ? __half2float(x[iy * kx + ix]) : 0.0f;
33
+ const float xi = ix < kx ? static_cast<float>(x[iy * kx + ix]) : 0.0f;
34
float amax = fabsf(xi);
35
float sum = xi;
36
@@ -51,14 +53,16 @@ static __global__ void quantize_q8_1(const half* __restrict__ x,
53
y[ib].ds.y = __float2half(sum);
54
}
55
54
-static void quantize_row_q8_1_cuda(const half* x, void* vy, const int kx,
56
+template <typename scalar_t>
57
+static void quantize_row_q8_1_cuda(const scalar_t* x, void* vy, const int kx,
58
const int ky, cudaStream_t stream) {
59
const int64_t kx_padded = (kx + 512 - 1) / 512 * 512;
60
const int block_num_x =
61
(kx_padded + CUDA_QUANTIZE_BLOCK_SIZE - 1) / CUDA_QUANTIZE_BLOCK_SIZE;
62
const dim3 num_blocks(block_num_x, ky, 1);
63
const dim3 block_size(CUDA_DEQUANTIZE_BLOCK_SIZE, 1, 1);
61
- quantize_q8_1<<<num_blocks, block_size, 0, stream>>>(x, vy, kx, kx_padded);
64
+ quantize_q8_1<scalar_t>
65
+ <<<num_blocks, block_size, 0, stream>>>(x, vy, kx, kx_padded);
66
}
67
68
torch::Tensor ggml_dequantize(torch::Tensor W, // quant weight
@@ -79,101 +83,112 @@ torch::Tensor ggml_mul_mat_vec_a8(torch::Tensor W, // quant weight
83
int col = X.sizes()[1];
84
const int padded = (col + 512 - 1) / 512 * 512;
85
const at::cuda::OptionalCUDAGuard device_guard(device_of(X));
82
- auto options =
83
- torch::TensorOptions().dtype(torch::kFloat16).device(W.device());
86
+ auto options = torch::TensorOptions().dtype(X.dtype()).device(W.device());
87
at::Tensor Y = torch::empty({1, row}, options);
88
cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream();
89
options = torch::TensorOptions().dtype(torch::kInt32).device(W.device());
90
at::Tensor quant_X = torch::empty({1, padded / 32 * 9}, options);
88
- quantize_row_q8_1_cuda((half*)X.data_ptr(), (void*)quant_X.data_ptr(), col, 1,
89
- stream);
90
- switch (type) {
91
- case 2:
92
- mul_mat_vec_q4_0_q8_1_cuda((void*)W.data_ptr(), (void*)quant_X.data_ptr(),
93
- (half*)Y.data_ptr(), col, row, stream);
94
- break;
95
- case 3:
96
- mul_mat_vec_q4_1_q8_1_cuda((void*)W.data_ptr(), (void*)quant_X.data_ptr(),
97
- (half*)Y.data_ptr(), col, row, stream);
98
- break;
99
- case 6:
100
- mul_mat_vec_q5_0_q8_1_cuda((void*)W.data_ptr(), (void*)quant_X.data_ptr(),
101
- (half*)Y.data_ptr(), col, row, stream);
102
- break;
103
- case 7:
104
- mul_mat_vec_q5_1_q8_1_cuda((void*)W.data_ptr(), (void*)quant_X.data_ptr(),
105
- (half*)Y.data_ptr(), col, row, stream);
106
- break;
107
- case 8:
108
- mul_mat_vec_q8_0_q8_1_cuda((void*)W.data_ptr(), (void*)quant_X.data_ptr(),
109
- (half*)Y.data_ptr(), col, row, stream);
110
- break;
111
- case 10:
112
- mul_mat_vec_q2_K_q8_1_cuda((void*)W.data_ptr(), (void*)quant_X.data_ptr(),
113
- (half*)Y.data_ptr(), col, row, stream);
114
- break;
115
- case 11:
116
- mul_mat_vec_q3_K_q8_1_cuda((void*)W.data_ptr(), (void*)quant_X.data_ptr(),
117
- (half*)Y.data_ptr(), col, row, stream);
118
- break;
119
- case 12:
120
- mul_mat_vec_q4_K_q8_1_cuda((void*)W.data_ptr(), (void*)quant_X.data_ptr(),
121
- (half*)Y.data_ptr(), col, row, stream);
122
- break;
123
- case 13:
124
- mul_mat_vec_q5_K_q8_1_cuda((void*)W.data_ptr(), (void*)quant_X.data_ptr(),
125
- (half*)Y.data_ptr(), col, row, stream);
126
- break;
127
- case 14:
128
- mul_mat_vec_q6_K_q8_1_cuda((void*)W.data_ptr(), (void*)quant_X.data_ptr(),
129
- (half*)Y.data_ptr(), col, row, stream);
130
- break;
131
- case 16:
132
- mul_mat_vec_iq2_xxs_q8_1_cuda((void*)W.data_ptr(),
133
- (void*)quant_X.data_ptr(),
134
- (half*)Y.data_ptr(), col, row, stream);
135
- break;
136
- case 17:
137
- mul_mat_vec_iq2_xs_q8_1_cuda((void*)W.data_ptr(),
138
- (void*)quant_X.data_ptr(),
139
- (half*)Y.data_ptr(), col, row, stream);
140
- break;
141
- case 18:
142
- mul_mat_vec_iq3_xxs_q8_1_cuda((void*)W.data_ptr(),
143
- (void*)quant_X.data_ptr(),
144
- (half*)Y.data_ptr(), col, row, stream);
145
- break;
146
- case 19:
147
- mul_mat_vec_iq1_s_q8_1_cuda((void*)W.data_ptr(),
148
- (void*)quant_X.data_ptr(),
149
- (half*)Y.data_ptr(), col, row, stream);
150
- break;
151
- case 20:
152
- mul_mat_vec_iq4_nl_q8_1_cuda((void*)W.data_ptr(),
153
- (void*)quant_X.data_ptr(),
154
- (half*)Y.data_ptr(), col, row, stream);
155
- break;
156
- case 21:
157
- mul_mat_vec_iq3_s_q8_1_cuda((void*)W.data_ptr(),
158
- (void*)quant_X.data_ptr(),
159
- (half*)Y.data_ptr(), col, row, stream);
160
- break;
161
- case 22:
162
- mul_mat_vec_iq2_s_q8_1_cuda((void*)W.data_ptr(),
163
- (void*)quant_X.data_ptr(),
164
- (half*)Y.data_ptr(), col, row, stream);
165
- break;
166
- case 23:
167
- mul_mat_vec_iq4_xs_q8_1_cuda((void*)W.data_ptr(),
168
- (void*)quant_X.data_ptr(),
169
- (half*)Y.data_ptr(), col, row, stream);
170
- break;
171
- case 29:
172
- mul_mat_vec_iq1_m_q8_1_cuda((void*)W.data_ptr(),
173
- (void*)quant_X.data_ptr(),
174
- (half*)Y.data_ptr(), col, row, stream);
175
- break;
176
- }
91
+ VLLM_DISPATCH_FLOATING_TYPES(X.scalar_type(), "ggml_mul_mat_vec_a8", [&] {
92
+ quantize_row_q8_1_cuda<scalar_t>((scalar_t*)X.data_ptr(),
93
+ (void*)quant_X.data_ptr(), col, 1, stream);
94
+ switch (type) {
95
+ case 2:
96
+ mul_mat_vec_q4_0_q8_1_cuda<scalar_t>(
97
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
98
+ (scalar_t*)Y.data_ptr(), col, row, stream);
99
+ break;
100
+ case 3:
101
+ mul_mat_vec_q4_1_q8_1_cuda<scalar_t>(
102
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
103
+ (scalar_t*)Y.data_ptr(), col, row, stream);
104
+ break;
105
+ case 6:
106
+ mul_mat_vec_q5_0_q8_1_cuda<scalar_t>(
107
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
108
+ (scalar_t*)Y.data_ptr(), col, row, stream);
109
+ break;
110
+ case 7:
111
+ mul_mat_vec_q5_1_q8_1_cuda<scalar_t>(
112
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
113
+ (scalar_t*)Y.data_ptr(), col, row, stream);
114
+ break;
115
+ case 8:
116
+ mul_mat_vec_q8_0_q8_1_cuda<scalar_t>(
117
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
118
+ (scalar_t*)Y.data_ptr(), col, row, stream);
119
+ break;
120
+ case 10:
121
+ mul_mat_vec_q2_K_q8_1_cuda<scalar_t>(
122
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
123
+ (scalar_t*)Y.data_ptr(), col, row, stream);
124
+ break;
125
+ case 11:
126
+ mul_mat_vec_q3_K_q8_1_cuda<scalar_t>(
127
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
128
+ (scalar_t*)Y.data_ptr(), col, row, stream);
129
+ break;
130
+ case 12:
131
+ mul_mat_vec_q4_K_q8_1_cuda<scalar_t>(
132
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
133
+ (scalar_t*)Y.data_ptr(), col, row, stream);
134
+ break;
135
+ case 13:
136
+ mul_mat_vec_q5_K_q8_1_cuda<scalar_t>(
137
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
138
+ (scalar_t*)Y.data_ptr(), col, row, stream);
139
+ break;
140
+ case 14:
141
+ mul_mat_vec_q6_K_q8_1_cuda<scalar_t>(
142
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
143
+ (scalar_t*)Y.data_ptr(), col, row, stream);
144
+ break;
145
+ case 16:
146
+ mul_mat_vec_iq2_xxs_q8_1_cuda<scalar_t>(
147
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
148
+ (scalar_t*)Y.data_ptr(), col, row, stream);
149
+ break;
150
+ case 17:
151
+ mul_mat_vec_iq2_xs_q8_1_cuda<scalar_t>(
152
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
153
+ (scalar_t*)Y.data_ptr(), col, row, stream);
154
+ break;
155
+ case 18:
156
+ mul_mat_vec_iq3_xxs_q8_1_cuda<scalar_t>(
157
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
158
+ (scalar_t*)Y.data_ptr(), col, row, stream);
159
+ break;
160
+ case 19:
161
+ mul_mat_vec_iq1_s_q8_1_cuda<scalar_t>(
162
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
163
+ (scalar_t*)Y.data_ptr(), col, row, stream);
164
+ break;
165
+ case 20:
166
+ mul_mat_vec_iq4_nl_q8_1_cuda<scalar_t>(
167
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
168
+ (scalar_t*)Y.data_ptr(), col, row, stream);
169
+ break;
170
+ case 21:
171
+ mul_mat_vec_iq3_s_q8_1_cuda<scalar_t>(
172
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
173
+ (scalar_t*)Y.data_ptr(), col, row, stream);
174
+ break;
175
+ case 22:
176
+ mul_mat_vec_iq2_s_q8_1_cuda<scalar_t>(
177
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
178
+ (scalar_t*)Y.data_ptr(), col, row, stream);
179
+ break;
180
+ case 23:
181
+ mul_mat_vec_iq4_xs_q8_1_cuda<scalar_t>(
182
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
183
+ (scalar_t*)Y.data_ptr(), col, row, stream);
184
+ break;
185
+ case 29:
186
+ mul_mat_vec_iq1_m_q8_1_cuda<scalar_t>(
187
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
188
+ (scalar_t*)Y.data_ptr(), col, row, stream);
189
+ break;
190
+ }
191
+ });
192
return Y;
193
}
194
@@ -184,66 +199,67 @@ torch::Tensor ggml_mul_mat_a8(torch::Tensor W, // quant weight
199
int padded = (col + 512 - 1) / 512 * 512;
200
int batch = X.sizes()[0];
201
const at::cuda::OptionalCUDAGuard device_guard(device_of(X));
187
- auto options =
188
- torch::TensorOptions().dtype(torch::kFloat16).device(W.device());
202
+ auto options = torch::TensorOptions().dtype(X.dtype()).device(W.device());
203
at::Tensor Y = torch::empty({batch, row}, options);
204
cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream();
205
options = torch::TensorOptions().dtype(torch::kInt32).device(W.device());
206
at::Tensor quant_X = torch::empty({batch, padded / 32 * 9}, options);
193
- quantize_row_q8_1_cuda((half*)X.data_ptr(), (void*)quant_X.data_ptr(), col,
194
- batch, stream);
195
-
196
- switch (type) {
197
- case 2:
198
- ggml_mul_mat_q4_0_q8_1_cuda(
199
- (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (half*)Y.data_ptr(),
200
- col, row, batch, padded, row, stream);
201
- break;
202
- case 3:
203
- ggml_mul_mat_q4_1_q8_1_cuda(
204
- (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (half*)Y.data_ptr(),
205
- col, row, batch, padded, row, stream);
206
- break;
207
- case 6:
208
- ggml_mul_mat_q5_0_q8_1_cuda(
209
- (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (half*)Y.data_ptr(),
210
- col, row, batch, padded, row, stream);
211
- break;
212
- case 7:
213
- ggml_mul_mat_q5_1_q8_1_cuda(
214
- (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (half*)Y.data_ptr(),
215
- col, row, batch, padded, row, stream);
216
- break;
217
- case 8:
218
- ggml_mul_mat_q8_0_q8_1_cuda(
219
- (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (half*)Y.data_ptr(),
220
- col, row, batch, padded, row, stream);
221
- break;
222
- case 10:
223
- ggml_mul_mat_q2_K_q8_1_cuda(
224
- (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (half*)Y.data_ptr(),
225
- col, row, batch, padded, row, stream);
226
- break;
227
- case 11:
228
- ggml_mul_mat_q3_K_q8_1_cuda(
229
- (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (half*)Y.data_ptr(),
230
- col, row, batch, padded, row, stream);
231
- break;
232
- case 12:
233
- ggml_mul_mat_q4_K_q8_1_cuda(
234
- (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (half*)Y.data_ptr(),
235
- col, row, batch, padded, row, stream);
236
- break;
237
- case 13:
238
- ggml_mul_mat_q5_K_q8_1_cuda(
239
- (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (half*)Y.data_ptr(),
240
- col, row, batch, padded, row, stream);
241
- break;
242
- case 14:
243
- ggml_mul_mat_q6_K_q8_1_cuda(
244
- (void*)W.data_ptr(), (void*)quant_X.data_ptr(), (half*)Y.data_ptr(),
245
- col, row, batch, padded, row, stream);
246
- break;
247
- }
207
+ VLLM_DISPATCH_FLOATING_TYPES(X.scalar_type(), "ggml_mul_mat_a8", [&] {
208
+ quantize_row_q8_1_cuda((scalar_t*)X.data_ptr(), (void*)quant_X.data_ptr(),
209
+ col, batch, stream);
210
+
211
+ switch (type) {
212
+ case 2:
213
+ ggml_mul_mat_q4_0_q8_1_cuda(
214
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
215
+ (scalar_t*)Y.data_ptr(), col, row, batch, padded, row, stream);
216
+ break;
217
+ case 3:
218
+ ggml_mul_mat_q4_1_q8_1_cuda(
219
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
220
+ (scalar_t*)Y.data_ptr(), col, row, batch, padded, row, stream);
221
+ break;
222
+ case 6:
223
+ ggml_mul_mat_q5_0_q8_1_cuda(
224
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
225
+ (scalar_t*)Y.data_ptr(), col, row, batch, padded, row, stream);
226
+ break;
227
+ case 7:
228
+ ggml_mul_mat_q5_1_q8_1_cuda(
229
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
230
+ (scalar_t*)Y.data_ptr(), col, row, batch, padded, row, stream);
231
+ break;
232
+ case 8:
233
+ ggml_mul_mat_q8_0_q8_1_cuda(
234
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
235
+ (scalar_t*)Y.data_ptr(), col, row, batch, padded, row, stream);
236
+ break;
237
+ case 10:
238
+ ggml_mul_mat_q2_K_q8_1_cuda(
239
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
240
+ (scalar_t*)Y.data_ptr(), col, row, batch, padded, row, stream);
241
+ break;
242
+ case 11:
243
+ ggml_mul_mat_q3_K_q8_1_cuda(
244
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
245
+ (scalar_t*)Y.data_ptr(), col, row, batch, padded, row, stream);
246
+ break;
247
+ case 12:
248
+ ggml_mul_mat_q4_K_q8_1_cuda(
249
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
250
+ (scalar_t*)Y.data_ptr(), col, row, batch, padded, row, stream);
251
+ break;
252
+ case 13:
253
+ ggml_mul_mat_q5_K_q8_1_cuda(
254
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
255
+ (scalar_t*)Y.data_ptr(), col, row, batch, padded, row, stream);
256
+ break;
257
+ case 14:
258
+ ggml_mul_mat_q6_K_q8_1_cuda(
259
+ (void*)W.data_ptr(), (void*)quant_X.data_ptr(),
260
+ (scalar_t*)Y.data_ptr(), col, row, batch, padded, row, stream);
261
+ break;
262
+ }
263
+ });
264
return Y;
265
}
csrc/quantization/gguf/mmq.cuh
+74
-64
@@ -1,8 +1,8 @@
1
// copied from https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/mmq.cu
2
-template <int qk, int qr, int qi, bool need_sum, typename block_q_t, int mmq_x, int mmq_y, int nwarps,
2
+template <typename scalar_t, int qk, int qr, int qi, bool need_sum, typename block_q_t, int mmq_x, int mmq_y, int nwarps,
3
allocate_tiles_cuda_t allocate_tiles, load_tiles_cuda_t load_tiles, int vdr, vec_dot_q_mul_mat_cuda_t vec_dot>
4
static __device__ __forceinline__ void mul_mat_q(
5
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
5
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
6
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
7
8
const block_q_t * x = (const block_q_t *) vx;
@@ -38,7 +38,7 @@ static __device__ __forceinline__ void mul_mat_q(
38
threadIdx.y, nrows_x-row_x_0-1, threadIdx.x, blocks_per_row_x);
39
40
#pragma unroll
41
- for (int ir = 0; ir < qr; ++ir) {
41
+ for (int ir = 0; ir < qr && ib0 + ir * blocks_per_warp/qr < blocks_per_row_x; ++ir) {
42
const int kqs = ir*WARP_SIZE_GGUF + threadIdx.x;
43
const int kbxd = kqs / QI8_1;
44
@@ -98,7 +98,7 @@ static __device__ __forceinline__ void mul_mat_q(
98
if (row_dst >= nrows_dst) {
99
continue;
100
}
101
- dst[col_dst*nrows_dst + row_dst] = __float2half(sum[i/WARP_SIZE_GGUF][j/nwarps]);
101
+ dst[col_dst*nrows_dst + row_dst] = sum[i/WARP_SIZE_GGUF][j/nwarps];
102
}
103
}
104
}
@@ -113,24 +113,25 @@ static __device__ __forceinline__ void mul_mat_q(
113
#define NWARPS_Q4_0 4
114
#endif
115
116
-template <bool need_check> static __global__ void
116
+template<typename scalar_t, bool need_check> static __global__ void
117
#if defined(USE_ROCM)
118
__launch_bounds__(WARP_SIZE_GGUF*NWARPS_Q4_0, 2)
119
#endif
120
mul_mat_q4_0(
121
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
121
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
122
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
123
const int mmq_x = MMQ_X_Q4_0;
124
const int mmq_y = MMQ_Y_Q4_0;
125
const int nwarps = NWARPS_Q4_0;
126
127
- mul_mat_q<QK4_0, QR4_0, QI4_0, true, block_q4_0, mmq_x, mmq_y, nwarps, allocate_tiles_q4_0<mmq_y>,
127
+ mul_mat_q<scalar_t, QK4_0, QR4_0, QI4_0, true, block_q4_0, mmq_x, mmq_y, nwarps, allocate_tiles_q4_0<mmq_y>,
128
load_tiles_q4_0<mmq_y, nwarps, need_check>, VDR_Q4_0_Q8_1_MMQ, vec_dot_q4_0_q8_1_mul_mat>
129
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
130
}
131
132
+template<typename scalar_t>
133
static void ggml_mul_mat_q4_0_q8_1_cuda(
133
- const void * vx, const void * vy, half * dst, const int ncols_x, const int nrows_x,
134
+ const void * vx, const void * vy, scalar_t * dst, const int ncols_x, const int nrows_x,
135
const int ncols_y, const int nrows_y, const int nrows_dst, cudaStream_t stream) {
136
137
int mmq_x = MMQ_X_Q4_0;
@@ -144,11 +145,11 @@ static void ggml_mul_mat_q4_0_q8_1_cuda(
145
146
if (nrows_x % mmq_y == 0) {
147
const bool need_check = false;
147
- mul_mat_q4_0<need_check><<<block_nums, block_dims, 0, stream>>>
148
+ mul_mat_q4_0<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
149
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
150
} else {
151
const bool need_check = true;
151
- mul_mat_q4_0<need_check><<<block_nums, block_dims, 0, stream>>>
152
+ mul_mat_q4_0<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
153
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
154
}
155
}
@@ -163,24 +164,25 @@ static void ggml_mul_mat_q4_0_q8_1_cuda(
164
#define NWARPS_Q4_1 4
165
#endif
166
166
-template <bool need_check> static __global__ void
167
+template<typename scalar_t, bool need_check> static __global__ void
168
#if defined(USE_ROCM)
169
__launch_bounds__(WARP_SIZE_GGUF*NWARPS_Q4_1, 2)
170
#endif
171
mul_mat_q4_1(
171
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
172
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
173
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
174
const int mmq_x = MMQ_X_Q4_1;
175
const int mmq_y = MMQ_Y_Q4_1;
176
const int nwarps = NWARPS_Q4_1;
177
177
- mul_mat_q<QK4_1, QR4_1, QI4_1, true, block_q4_1, mmq_x, mmq_y, nwarps, allocate_tiles_q4_1<mmq_y>,
178
+ mul_mat_q<scalar_t, QK4_1, QR4_1, QI4_1, true, block_q4_1, mmq_x, mmq_y, nwarps, allocate_tiles_q4_1<mmq_y>,
179
load_tiles_q4_1<mmq_y, nwarps, need_check>, VDR_Q4_1_Q8_1_MMQ, vec_dot_q4_1_q8_1_mul_mat>
180
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
181
}
182
183
+template<typename scalar_t>
184
static void ggml_mul_mat_q4_1_q8_1_cuda(
183
- const void * vx, const void * vy, half * dst, const int ncols_x, const int nrows_x,
185
+ const void * vx, const void * vy, scalar_t * dst, const int ncols_x, const int nrows_x,
186
const int ncols_y, const int nrows_y, const int nrows_dst, cudaStream_t stream) {
187
188
int mmq_x = MMQ_X_Q4_1;
@@ -194,11 +196,11 @@ static void ggml_mul_mat_q4_1_q8_1_cuda(
196
197
if (nrows_x % mmq_y == 0) {
198
const bool need_check = false;
197
- mul_mat_q4_1<need_check><<<block_nums, block_dims, 0, stream>>>
199
+ mul_mat_q4_1<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
200
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
201
} else {
202
const bool need_check = true;
201
- mul_mat_q4_1<need_check><<<block_nums, block_dims, 0, stream>>>
203
+ mul_mat_q4_1<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
204
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
205
}
206
}
@@ -213,24 +215,25 @@ static void ggml_mul_mat_q4_1_q8_1_cuda(
215
#define NWARPS_Q5_0 4
216
#endif
217
216
-template <bool need_check> static __global__ void
218
+template<typename scalar_t, bool need_check> static __global__ void
219
#if defined(USE_ROCM)
220
__launch_bounds__(WARP_SIZE_GGUF*NWARPS_Q5_0, 2)
221
#endif
222
mul_mat_q5_0(
221
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
223
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
224
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
225
const int mmq_x = MMQ_X_Q5_0;
226
const int mmq_y = MMQ_Y_Q5_0;
227
const int nwarps = NWARPS_Q5_0;
228
227
- mul_mat_q<QK5_0, QR5_0, QI5_0, false, block_q5_0, mmq_x, mmq_y, nwarps, allocate_tiles_q5_0<mmq_y>,
229
+ mul_mat_q<scalar_t, QK5_0, QR5_0, QI5_0, false, block_q5_0, mmq_x, mmq_y, nwarps, allocate_tiles_q5_0<mmq_y>,
230
load_tiles_q5_0<mmq_y, nwarps, need_check>, VDR_Q5_0_Q8_1_MMQ, vec_dot_q5_0_q8_1_mul_mat>
231
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
232
}
233
234
+template<typename scalar_t>
235
static void ggml_mul_mat_q5_0_q8_1_cuda(
233
- const void * vx, const void * vy, half * dst, const int ncols_x, const int nrows_x,
236
+ const void * vx, const void * vy, scalar_t * dst, const int ncols_x, const int nrows_x,
237
const int ncols_y, const int nrows_y, const int nrows_dst, cudaStream_t stream) {
238
239
const int mmq_x = MMQ_X_Q5_0;
@@ -244,11 +247,11 @@ static void ggml_mul_mat_q5_0_q8_1_cuda(
247
248
if (nrows_x % mmq_y == 0) {
249
const bool need_check = false;
247
- mul_mat_q5_0<need_check><<<block_nums, block_dims, 0, stream>>>
250
+ mul_mat_q5_0<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
251
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
252
} else {
253
const bool need_check = true;
251
- mul_mat_q5_0<need_check><<<block_nums, block_dims, 0, stream>>>
254
+ mul_mat_q5_0<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
255
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
256
}
257
}
@@ -263,24 +266,25 @@ static void ggml_mul_mat_q5_0_q8_1_cuda(
266
#define NWARPS_Q5_1 4
267
#endif
268
266
-template <bool need_check> static __global__ void
269
+template<typename scalar_t, bool need_check> static __global__ void
270
#if defined(USE_ROCM)
271
__launch_bounds__(WARP_SIZE_GGUF*NWARPS_Q5_1, 2)
272
#endif
273
mul_mat_q5_1(
271
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
274
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
275
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
276
const int mmq_x = MMQ_X_Q5_1;
277
const int mmq_y = MMQ_Y_Q5_1;
278
const int nwarps = NWARPS_Q5_1;
279
277
- mul_mat_q<QK5_1, QR5_1, QI5_1, true, block_q5_1, mmq_x, mmq_y, nwarps, allocate_tiles_q5_1<mmq_y>,
280
+ mul_mat_q<scalar_t, QK5_1, QR5_1, QI5_1, true, block_q5_1, mmq_x, mmq_y, nwarps, allocate_tiles_q5_1<mmq_y>,
281
load_tiles_q5_1<mmq_y, nwarps, need_check>, VDR_Q5_1_Q8_1_MMQ, vec_dot_q5_1_q8_1_mul_mat>
282
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
283
}
284
285
+template<typename scalar_t>
286
static void ggml_mul_mat_q5_1_q8_1_cuda(
283
- const void * vx, const void * vy, half * dst, const int ncols_x, const int nrows_x,
287
+ const void * vx, const void * vy, scalar_t * dst, const int ncols_x, const int nrows_x,
288
const int ncols_y, const int nrows_y, const int nrows_dst, cudaStream_t stream) {
289
const int mmq_x = MMQ_X_Q5_1;
290
const int mmq_y = MMQ_Y_Q5_1;
@@ -293,11 +297,11 @@ static void ggml_mul_mat_q5_1_q8_1_cuda(
297
298
if (nrows_x % mmq_y == 0) {
299
const bool need_check = false;
296
- mul_mat_q5_1<need_check><<<block_nums, block_dims, 0, stream>>>
300
+ mul_mat_q5_1<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
301
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
302
} else {
303
const bool need_check = true;
300
- mul_mat_q5_1<need_check><<<block_nums, block_dims, 0, stream>>>
304
+ mul_mat_q5_1<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
305
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
306
}
307
}
@@ -312,24 +316,25 @@ static void ggml_mul_mat_q5_1_q8_1_cuda(
316
#define NWARPS_Q8_0 4
317
#endif
318
315
-template <bool need_check> static __global__ void
319
+template<typename scalar_t, bool need_check> static __global__ void
320
#if defined(USE_ROCM)
321
__launch_bounds__(WARP_SIZE_GGUF*NWARPS_Q8_0, 2)
322
#endif
323
mul_mat_q8_0(
320
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
324
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
325
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
326
const int mmq_x = MMQ_X_Q8_0;
327
const int mmq_y = MMQ_Y_Q8_0;
328
const int nwarps = NWARPS_Q8_0;
329
326
- mul_mat_q<QK8_0, QR8_0, QI8_0, false, block_q8_0, mmq_x, mmq_y, nwarps, allocate_tiles_q8_0<mmq_y>,
330
+ mul_mat_q<scalar_t, QK8_0, QR8_0, QI8_0, false, block_q8_0, mmq_x, mmq_y, nwarps, allocate_tiles_q8_0<mmq_y>,
331
load_tiles_q8_0<mmq_y, nwarps, need_check>, VDR_Q8_0_Q8_1_MMQ, vec_dot_q8_0_q8_1_mul_mat>
332
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
333
}
334
335
+template<typename scalar_t>
336
static void ggml_mul_mat_q8_0_q8_1_cuda(
332
- const void * vx, const void * vy, half * dst, const int ncols_x, const int nrows_x,
337
+ const void * vx, const void * vy, scalar_t * dst, const int ncols_x, const int nrows_x,
338
const int ncols_y, const int nrows_y, const int nrows_dst, cudaStream_t stream) {
339
const int mmq_x = MMQ_X_Q8_0;
340
const int mmq_y = MMQ_Y_Q8_0;
@@ -342,11 +347,11 @@ static void ggml_mul_mat_q8_0_q8_1_cuda(
347
348
if (nrows_x % mmq_y == 0) {
349
const bool need_check = false;
345
- mul_mat_q8_0<need_check><<<block_nums, block_dims, 0, stream>>>
350
+ mul_mat_q8_0<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
351
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
352
} else {
353
const bool need_check = true;
349
- mul_mat_q8_0<need_check><<<block_nums, block_dims, 0, stream>>>
354
+ mul_mat_q8_0<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
355
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
356
}
357
}
@@ -361,24 +366,25 @@ static void ggml_mul_mat_q8_0_q8_1_cuda(
366
#define NWARPS_Q2_K 4
367
#endif
368
364
-template <bool need_check> static __global__ void
369
+template<typename scalar_t, bool need_check> static __global__ void
370
#if defined(USE_ROCM)
371
__launch_bounds__(WARP_SIZE_GGUF*NWARPS_Q2_K, 2)
372
#endif
373
mul_mat_q2_K(
369
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
374
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
375
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
376
const int mmq_x = MMQ_X_Q2_K;
377
const int mmq_y = MMQ_Y_Q2_K;
378
const int nwarps = NWARPS_Q2_K;
379
375
- mul_mat_q<QK_K, QR2_K, QI2_K, false, block_q2_K, mmq_x, mmq_y, nwarps, allocate_tiles_q2_K<mmq_y>,
380
+ mul_mat_q<scalar_t, QK_K, QR2_K, QI2_K, false, block_q2_K, mmq_x, mmq_y, nwarps, allocate_tiles_q2_K<mmq_y>,
381
load_tiles_q2_K<mmq_y, nwarps, need_check>, VDR_Q2_K_Q8_1_MMQ, vec_dot_q2_K_q8_1_mul_mat>
382
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
383
}
384
385
+template<typename scalar_t>
386
static void ggml_mul_mat_q2_K_q8_1_cuda(
381
- const void * vx, const void * vy, half * dst, const int ncols_x, const int nrows_x,
387
+ const void * vx, const void * vy, scalar_t * dst, const int ncols_x, const int nrows_x,
388
const int ncols_y, const int nrows_y, const int nrows_dst, cudaStream_t stream) {
389
const int mmq_x = MMQ_X_Q2_K;
390
const int mmq_y = MMQ_Y_Q2_K;
@@ -391,11 +397,11 @@ static void ggml_mul_mat_q2_K_q8_1_cuda(
397
398
if (nrows_x % mmq_y == 0) {
399
const bool need_check = false;
394
- mul_mat_q2_K<need_check><<<block_nums, block_dims, 0, stream>>>
400
+ mul_mat_q2_K<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
401
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
402
} else {
403
const bool need_check = true;
398
- mul_mat_q2_K<need_check><<<block_nums, block_dims, 0, stream>>>
404
+ mul_mat_q2_K<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
405
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
406
}
407
}
@@ -410,25 +416,26 @@ static void ggml_mul_mat_q2_K_q8_1_cuda(
416
#define NWARPS_Q3_K 4
417
#endif
418
413
-template <bool need_check> static __global__ void
419
+template<typename scalar_t, bool need_check> static __global__ void
420
#if defined(USE_ROCM)
421
__launch_bounds__(WARP_SIZE_GGUF*NWARPS_Q3_K, 2)
422
#endif
423
mul_mat_q3_K(
418
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
424
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
425
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
426
427
const int mmq_x = MMQ_X_Q3_K;
428
const int mmq_y = MMQ_Y_Q3_K;
429
const int nwarps = NWARPS_Q3_K;
430
425
- mul_mat_q<QK_K, QR3_K, QI3_K, false, block_q3_K, mmq_x, mmq_y, nwarps, allocate_tiles_q3_K<mmq_y>,
431
+ mul_mat_q<scalar_t, QK_K, QR3_K, QI3_K, false, block_q3_K, mmq_x, mmq_y, nwarps, allocate_tiles_q3_K<mmq_y>,
432
load_tiles_q3_K<mmq_y, nwarps, need_check>, VDR_Q3_K_Q8_1_MMQ, vec_dot_q3_K_q8_1_mul_mat>
433
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
434
}
435
436
+template<typename scalar_t>
437
static void ggml_mul_mat_q3_K_q8_1_cuda(
431
- const void * vx, const void * vy, half * dst, const int ncols_x, const int nrows_x,
438
+ const void * vx, const void * vy, scalar_t * dst, const int ncols_x, const int nrows_x,
439
const int ncols_y, const int nrows_y, const int nrows_dst, cudaStream_t stream) {
440
441
const int mmq_x = MMQ_X_Q3_K;
@@ -442,11 +449,11 @@ static void ggml_mul_mat_q3_K_q8_1_cuda(
449
450
if (nrows_x % mmq_y == 0) {
451
const bool need_check = false;
445
- mul_mat_q3_K<need_check><<<block_nums, block_dims, 0, stream>>>
452
+ mul_mat_q3_K<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
453
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
454
} else {
455
const bool need_check = true;
449
- mul_mat_q3_K<need_check><<<block_nums, block_dims, 0, stream>>>
456
+ mul_mat_q3_K<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
457
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
458
}
459
}
@@ -461,24 +468,25 @@ static void ggml_mul_mat_q3_K_q8_1_cuda(
468
#define NWARPS_Q4_K 4
469
#endif
470
464
-template <bool need_check> static __global__ void
471
+template<typename scalar_t, bool need_check> static __global__ void
472
#if defined(USE_ROCM)
473
__launch_bounds__(WARP_SIZE_GGUF*NWARPS_Q4_K, 2)
474
#endif
475
mul_mat_q4_K(
469
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
476
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
477
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
478
const int mmq_x = MMQ_X_Q4_K;
479
const int mmq_y = MMQ_Y_Q4_K;
480
const int nwarps = NWARPS_Q4_K;
481
475
- mul_mat_q<QK_K, QR4_K, QI4_K, true, block_q4_K, mmq_x, mmq_y, nwarps, allocate_tiles_q4_K<mmq_y>,
482
+ mul_mat_q<scalar_t, QK_K, QR4_K, QI4_K, true, block_q4_K, mmq_x, mmq_y, nwarps, allocate_tiles_q4_K<mmq_y>,
483
load_tiles_q4_K<mmq_y, nwarps, need_check>, VDR_Q4_K_Q8_1_MMQ, vec_dot_q4_K_q8_1_mul_mat>
484
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
485
}
486
487
+template<typename scalar_t>
488
static void ggml_mul_mat_q4_K_q8_1_cuda(
481
- const void * vx, const void * vy, half * dst, const int ncols_x, const int nrows_x,
489
+ const void * vx, const void * vy, scalar_t * dst, const int ncols_x, const int nrows_x,
490
const int ncols_y, const int nrows_y, const int nrows_dst, cudaStream_t stream) {
491
const int mmq_x = MMQ_X_Q4_K;
492
const int mmq_y = MMQ_Y_Q4_K;
@@ -491,11 +499,11 @@ static void ggml_mul_mat_q4_K_q8_1_cuda(
499
500
if (nrows_x % mmq_y == 0) {
501
const bool need_check = false;
494
- mul_mat_q4_K<need_check><<<block_nums, block_dims, 0, stream>>>
502
+ mul_mat_q4_K<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
503
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
504
} else {
505
const bool need_check = true;
498
- mul_mat_q4_K<need_check><<<block_nums, block_dims, 0, stream>>>
506
+ mul_mat_q4_K<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
507
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
508
}
509
}
@@ -510,24 +518,25 @@ static void ggml_mul_mat_q4_K_q8_1_cuda(
518
#define NWARPS_Q5_K 4
519
#endif
520
513
-template <bool need_check> static __global__ void
521
+template<typename scalar_t, bool need_check> static __global__ void
522
#if defined(USE_ROCM)
523
__launch_bounds__(WARP_SIZE_GGUF*NWARPS_Q5_K, 2)
524
#endif
525
mul_mat_q5_K(
518
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
526
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
527
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
528
const int mmq_x = MMQ_X_Q5_K;
529
const int mmq_y = MMQ_Y_Q5_K;
530
const int nwarps = NWARPS_Q5_K;
531
524
- mul_mat_q<QK_K, QR5_K, QI5_K, true, block_q5_K, mmq_x, mmq_y, nwarps, allocate_tiles_q5_K<mmq_y>,
532
+ mul_mat_q<scalar_t, QK_K, QR5_K, QI5_K, true, block_q5_K, mmq_x, mmq_y, nwarps, allocate_tiles_q5_K<mmq_y>,
533
load_tiles_q5_K<mmq_y, nwarps, need_check>, VDR_Q5_K_Q8_1_MMQ, vec_dot_q5_K_q8_1_mul_mat>
534
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
535
}
536
537
+template<typename scalar_t>
538
static void ggml_mul_mat_q5_K_q8_1_cuda(
530
- const void * vx, const void * vy, half * dst, const int ncols_x, const int nrows_x,
539
+ const void * vx, const void * vy, scalar_t * dst, const int ncols_x, const int nrows_x,
540
const int ncols_y, const int nrows_y, const int nrows_dst, cudaStream_t stream) {
541
542
const int mmq_x = MMQ_X_Q5_K;
@@ -541,11 +550,11 @@ static void ggml_mul_mat_q5_K_q8_1_cuda(
550
551
if (nrows_x % mmq_y == 0) {
552
const bool need_check = false;
544
- mul_mat_q5_K<need_check><<<block_nums, block_dims, 0, stream>>>
553
+ mul_mat_q5_K<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
554
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
555
} else {
556
const bool need_check = true;
548
- mul_mat_q5_K<need_check><<<block_nums, block_dims, 0, stream>>>
557
+ mul_mat_q5_K<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
558
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
559
}
560
}
@@ -560,24 +569,25 @@ static void ggml_mul_mat_q5_K_q8_1_cuda(
569
#define NWARPS_Q6_K 4
570
#endif
571
563
-template <bool need_check> static __global__ void
572
+template<typename scalar_t, bool need_check> static __global__ void
573
#if defined(USE_ROCM)
574
__launch_bounds__(WARP_SIZE_GGUF*NWARPS_Q6_K, 2)
575
#endif
576
mul_mat_q6_K(
568
- const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst,
577
+ const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst,
578
const int ncols_x, const int nrows_x, const int ncols_y, const int nrows_y, const int nrows_dst) {
579
const int mmq_x = MMQ_X_Q6_K;
580
const int mmq_y = MMQ_Y_Q6_K;
581
const int nwarps = NWARPS_Q6_K;
582
574
- mul_mat_q<QK_K, QR6_K, QI6_K, false, block_q6_K, mmq_x, mmq_y, nwarps, allocate_tiles_q6_K<mmq_y>,
583
+ mul_mat_q<scalar_t, QK_K, QR6_K, QI6_K, false, block_q6_K, mmq_x, mmq_y, nwarps, allocate_tiles_q6_K<mmq_y>,
584
load_tiles_q6_K<mmq_y, nwarps, need_check>, VDR_Q6_K_Q8_1_MMQ, vec_dot_q6_K_q8_1_mul_mat>
585
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
586
}
587
588
+template<typename scalar_t>
589
static void ggml_mul_mat_q6_K_q8_1_cuda(
580
- const void * vx, const void * vy, half * dst, const int ncols_x, const int nrows_x,
590
+ const void * vx, const void * vy, scalar_t * dst, const int ncols_x, const int nrows_x,
591
const int ncols_y, const int nrows_y, const int nrows_dst, cudaStream_t stream) {
592
const int mmq_x = MMQ_X_Q6_K;
593
const int mmq_y = MMQ_Y_Q6_K;
@@ -590,11 +600,11 @@ static void ggml_mul_mat_q6_K_q8_1_cuda(
600
601
if (nrows_x % mmq_y == 0) {
602
const bool need_check = false;
593
- mul_mat_q6_K<need_check><<<block_nums, block_dims, 0, stream>>>
603
+ mul_mat_q6_K<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
604
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
605
} else {
606
const bool need_check = true;
597
- mul_mat_q6_K<need_check><<<block_nums, block_dims, 0, stream>>>
607
+ mul_mat_q6_K<scalar_t, need_check><<<block_nums, block_dims, 0, stream>>>
608
(vx, vy, dst, ncols_x, nrows_x, ncols_y, nrows_y, nrows_dst);
609
}
610
}
csrc/quantization/gguf/mmvq.cuh
+60
-41
@@ -1,6 +1,6 @@
1
// copied and adapted from https://github.com/ggerganov/llama.cpp/blob/b2899/ggml-cuda/mmvq.cu
2
-template <int qk, int qi, typename block_q_t, int vdr, vec_dot_q_cuda_t vec_dot_q_cuda>
3
-static __global__ void mul_mat_vec_q(const void * __restrict__ vx, const void * __restrict__ vy, half * __restrict__ dst, const int ncols, const int nrows) {
2
+template <typename scalar_t, int qk, int qi, typename block_q_t, int vdr, vec_dot_q_cuda_t vec_dot_q_cuda>
3
+static __global__ void mul_mat_vec_q(const void * __restrict__ vx, const void * __restrict__ vy, scalar_t * __restrict__ dst, const int ncols, const int nrows) {
4
const int row = blockIdx.x*blockDim.y + threadIdx.y;
5
6
if (row >= nrows) {
@@ -33,158 +33,177 @@ static __global__ void mul_mat_vec_q(const void * __restrict__ vx, const void *
33
}
34
35
if (threadIdx.x == 0) {
36
- dst[row] = __float2half(tmp);
36
+ dst[row] = tmp;
37
}
38
}
39
40
-static void mul_mat_vec_q4_0_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
40
+template<typename scalar_t>
41
+static void mul_mat_vec_q4_0_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
42
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
43
const dim3 block_nums(block_num_y, 1, 1);
44
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
44
- mul_mat_vec_q<QK4_0, QI4_0, block_q4_0, VDR_Q4_0_Q8_1_MMVQ, vec_dot_q4_0_q8_1>
45
+ mul_mat_vec_q<scalar_t, QK4_0, QI4_0, block_q4_0, VDR_Q4_0_Q8_1_MMVQ, vec_dot_q4_0_q8_1>
46
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
47
}
48
48
-static void mul_mat_vec_q4_1_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
49
+template<typename scalar_t>
50
+static void mul_mat_vec_q4_1_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
51
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
52
const dim3 block_nums(block_num_y, 1, 1);
53
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
52
- mul_mat_vec_q<QK4_0, QI4_1, block_q4_1, VDR_Q4_1_Q8_1_MMVQ, vec_dot_q4_1_q8_1>
54
+ mul_mat_vec_q<scalar_t, QK4_0, QI4_1, block_q4_1, VDR_Q4_1_Q8_1_MMVQ, vec_dot_q4_1_q8_1>
55
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
56
}
57
56
-static void mul_mat_vec_q5_0_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
58
+template<typename scalar_t>
59
+static void mul_mat_vec_q5_0_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
60
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
61
const dim3 block_nums(block_num_y, 1, 1);
62
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
60
- mul_mat_vec_q<QK5_0, QI5_0, block_q5_0, VDR_Q5_0_Q8_1_MMVQ, vec_dot_q5_0_q8_1>
63
+ mul_mat_vec_q<scalar_t, QK5_0, QI5_0, block_q5_0, VDR_Q5_0_Q8_1_MMVQ, vec_dot_q5_0_q8_1>
64
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
65
}
66
64
-static void mul_mat_vec_q5_1_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
67
+template<typename scalar_t>
68
+static void mul_mat_vec_q5_1_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
69
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
70
const dim3 block_nums(block_num_y, 1, 1);
71
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
68
- mul_mat_vec_q<QK5_1, QI5_1, block_q5_1, VDR_Q5_1_Q8_1_MMVQ, vec_dot_q5_1_q8_1>
72
+ mul_mat_vec_q<scalar_t, QK5_1, QI5_1, block_q5_1, VDR_Q5_1_Q8_1_MMVQ, vec_dot_q5_1_q8_1>
73
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
74
}
75
72
-static void mul_mat_vec_q8_0_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
76
+template<typename scalar_t>
77
+static void mul_mat_vec_q8_0_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
78
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
79
const dim3 block_nums(block_num_y, 1, 1);
80
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
76
- mul_mat_vec_q<QK8_0, QI8_0, block_q8_0, VDR_Q8_0_Q8_1_MMVQ, vec_dot_q8_0_q8_1>
81
+ mul_mat_vec_q<scalar_t, QK8_0, QI8_0, block_q8_0, VDR_Q8_0_Q8_1_MMVQ, vec_dot_q8_0_q8_1>
82
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
83
}
84
80
-static void mul_mat_vec_q2_K_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
85
+template<typename scalar_t>
86
+static void mul_mat_vec_q2_K_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
87
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
88
const dim3 block_nums(block_num_y, 1, 1);
89
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
84
- mul_mat_vec_q<QK_K, QI2_K, block_q2_K, VDR_Q2_K_Q8_1_MMVQ, vec_dot_q2_K_q8_1>
90
+ mul_mat_vec_q<scalar_t, QK_K, QI2_K, block_q2_K, VDR_Q2_K_Q8_1_MMVQ, vec_dot_q2_K_q8_1>
91
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
92
}
93
88
-static void mul_mat_vec_q3_K_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
94
+template<typename scalar_t>
95
+static void mul_mat_vec_q3_K_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
96
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
97
const dim3 block_nums(block_num_y, 1, 1);
98
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
92
- mul_mat_vec_q<QK_K, QI3_K, block_q3_K, VDR_Q3_K_Q8_1_MMVQ, vec_dot_q3_K_q8_1>
99
+ mul_mat_vec_q<scalar_t, QK_K, QI3_K, block_q3_K, VDR_Q3_K_Q8_1_MMVQ, vec_dot_q3_K_q8_1>
100
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
101
}
102
96
-static void mul_mat_vec_q4_K_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
103
+template<typename scalar_t>
104
+static void mul_mat_vec_q4_K_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
105
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
106
const dim3 block_nums(block_num_y, 1, 1);
107
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
100
- mul_mat_vec_q<QK_K, QI4_K, block_q4_K, VDR_Q4_K_Q8_1_MMVQ, vec_dot_q4_K_q8_1>
108
+ mul_mat_vec_q<scalar_t, QK_K, QI4_K, block_q4_K, VDR_Q4_K_Q8_1_MMVQ, vec_dot_q4_K_q8_1>
109
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
110
}
111
104
-static void mul_mat_vec_q5_K_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
112
+template<typename scalar_t>
113
+static void mul_mat_vec_q5_K_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
114
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
115
const dim3 block_nums(block_num_y, 1, 1);
116
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
108
- mul_mat_vec_q<QK_K, QI5_K, block_q5_K, VDR_Q5_K_Q8_1_MMVQ, vec_dot_q5_K_q8_1>
117
+ mul_mat_vec_q<scalar_t, QK_K, QI5_K, block_q5_K, VDR_Q5_K_Q8_1_MMVQ, vec_dot_q5_K_q8_1>
118
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
119
}
120
112
-static void mul_mat_vec_q6_K_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
121
+template<typename scalar_t>
122
+static void mul_mat_vec_q6_K_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
123
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
124
const dim3 block_nums(block_num_y, 1, 1);
125
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
116
- mul_mat_vec_q<QK_K, QI6_K, block_q6_K, VDR_Q6_K_Q8_1_MMVQ, vec_dot_q6_K_q8_1>
126
+ mul_mat_vec_q<scalar_t, QK_K, QI6_K, block_q6_K, VDR_Q6_K_Q8_1_MMVQ, vec_dot_q6_K_q8_1>
127
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
128
}
129
120
-static void mul_mat_vec_iq2_xxs_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
130
+template<typename scalar_t>
131
+static void mul_mat_vec_iq2_xxs_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
132
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
133
const dim3 block_nums(block_num_y, 1, 1);
134
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
124
- mul_mat_vec_q<QK_K, QI2_XXS, block_iq2_xxs, 1, vec_dot_iq2_xxs_q8_1>
135
+ mul_mat_vec_q<scalar_t, QK_K, QI2_XXS, block_iq2_xxs, 1, vec_dot_iq2_xxs_q8_1>
136
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
137
}
138
128
-static void mul_mat_vec_iq2_xs_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
139
+template<typename scalar_t>
140
+static void mul_mat_vec_iq2_xs_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
141
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
142
const dim3 block_nums(block_num_y, 1, 1);
143
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
132
- mul_mat_vec_q<QK_K, QI2_XS, block_iq2_xs, 1, vec_dot_iq2_xs_q8_1>
144
+ mul_mat_vec_q<scalar_t, QK_K, QI2_XS, block_iq2_xs, 1, vec_dot_iq2_xs_q8_1>
145
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
146
}
147
136
-static void mul_mat_vec_iq2_s_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
148
+template<typename scalar_t>
149
+static void mul_mat_vec_iq2_s_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
150
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
151
const dim3 block_nums(block_num_y, 1, 1);
152
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
140
- mul_mat_vec_q<QK_K, QI2_S, block_iq2_s, 1, vec_dot_iq2_s_q8_1>
153
+ mul_mat_vec_q<scalar_t, QK_K, QI2_S, block_iq2_s, 1, vec_dot_iq2_s_q8_1>
154
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
155
}
156
144
-static void mul_mat_vec_iq3_xxs_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
157
+template<typename scalar_t>
158
+static void mul_mat_vec_iq3_xxs_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
159
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
160
const dim3 block_nums(block_num_y, 1, 1);
161
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
148
- mul_mat_vec_q<QK_K, QI3_XXS, block_iq3_xxs, 1, vec_dot_iq3_xxs_q8_1>
162
+ mul_mat_vec_q<scalar_t, QK_K, QI3_XXS, block_iq3_xxs, 1, vec_dot_iq3_xxs_q8_1>
163
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
164
}
165
152
-static void mul_mat_vec_iq1_s_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
166
+template<typename scalar_t>
167
+static void mul_mat_vec_iq1_s_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
168
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
169
const dim3 block_nums(block_num_y, 1, 1);
170
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
156
- mul_mat_vec_q<QK_K, QI1_S, block_iq1_s, 1, vec_dot_iq1_s_q8_1>
171
+ mul_mat_vec_q<scalar_t, QK_K, QI1_S, block_iq1_s, 1, vec_dot_iq1_s_q8_1>
172
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
173
}
174
160
-static void mul_mat_vec_iq1_m_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
175
+template<typename scalar_t>
176
+static void mul_mat_vec_iq1_m_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
177
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
178
const dim3 block_nums(block_num_y, 1, 1);
179
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
164
- mul_mat_vec_q<QK_K, QI1_M, block_iq1_m, 1, vec_dot_iq1_m_q8_1>
180
+ mul_mat_vec_q<scalar_t, QK_K, QI1_M, block_iq1_m, 1, vec_dot_iq1_m_q8_1>
181
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
182
}
183
168
-static void mul_mat_vec_iq4_nl_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
184
+template<typename scalar_t>
185
+static void mul_mat_vec_iq4_nl_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
186
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
187
const dim3 block_nums(block_num_y, 1, 1);
188
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
172
- mul_mat_vec_q<QK4_NL, QI4_NL, block_iq4_nl, VDR_Q4_0_Q8_1_MMVQ, vec_dot_iq4_nl_q8_1>
189
+ mul_mat_vec_q<scalar_t, QK4_NL, QI4_NL, block_iq4_nl, VDR_Q4_0_Q8_1_MMVQ, vec_dot_iq4_nl_q8_1>
190
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
191
}
192
176
-static void mul_mat_vec_iq4_xs_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
193
+template<typename scalar_t>
194
+static void mul_mat_vec_iq4_xs_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
195
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
196
const dim3 block_nums(block_num_y, 1, 1);
197
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
180
- mul_mat_vec_q<QK_K, QI4_XS, block_iq4_xs, 1, vec_dot_iq4_xs_q8_1>
198
+ mul_mat_vec_q<scalar_t, QK_K, QI4_XS, block_iq4_xs, 1, vec_dot_iq4_xs_q8_1>
199
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
200
}
201
184
-static void mul_mat_vec_iq3_s_q8_1_cuda(const void * vx, const void * vy, half * dst, const int ncols, const int nrows, cudaStream_t stream) {
202
+template<typename scalar_t>
203
+static void mul_mat_vec_iq3_s_q8_1_cuda(const void * vx, const void * vy, scalar_t * dst, const int ncols, const int nrows, cudaStream_t stream) {
204
const int block_num_y = (nrows + GGML_CUDA_MMV_Y - 1) / GGML_CUDA_MMV_Y;
205
const dim3 block_nums(block_num_y, 1, 1);
206
const dim3 block_dims(WARP_SIZE, GGML_CUDA_MMV_Y, 1);
188
- mul_mat_vec_q<QK_K, QI3_XS, block_iq3_s, 1, vec_dot_iq3_s_q8_1>
207
+ mul_mat_vec_q<scalar_t, QK_K, QI3_XS, block_iq3_s, 1, vec_dot_iq3_s_q8_1>
208
<<<block_nums, block_dims, 0, stream>>>(vx, vy, dst, ncols, nrows);
209
}
tests/kernels/test_gguf.py
+12
-6
@@ -22,7 +22,7 @@ def get_gguf_sample_tensors(
22
return GGUFReader(sample_file).tensors
23
24
25
-DTYPES = [torch.half]
25
+DTYPES = [torch.half, torch.bfloat16, torch.float32]
26
# Hidden_size for testing, must match the sample file in HF repo,
27
# we have `hidden_size = 256, 1024` for test in HF repo currently.
28
HIDDEN_SIZES = [256, 1024]
@@ -52,7 +52,7 @@ QUANT_TYPES = [
52
53
54
@pytest.mark.parametrize("hidden_size", HIDDEN_SIZES)
55
-@pytest.mark.parametrize("dtype", DTYPES)
55
+@pytest.mark.parametrize("dtype", [torch.half])
56
@pytest.mark.parametrize("quant_type", QUANT_TYPES)
57
@torch.inference_mode()
58
def test_dequantize(hidden_size: int, dtype: torch.dtype,
@@ -122,7 +122,13 @@ def test_mmq(num_tokens: int, hidden_size: int, dtype: torch.dtype,
122
ref_output = x @ weight.T
123
124
qweight = torch.tensor(tensor.data, device="cuda")
125
- output = ops.ggml_mul_mat_a8(qweight, x, quant_type,
126
- qweight.shape[0]).to(dtype)
127
-
128
- torch.testing.assert_close(output, ref_output, atol=1, rtol=1e-1)
125
+ output = ops.ggml_mul_mat_a8(qweight, x, quant_type, qweight.shape[0])
126
+ atols = {torch.half: 1, torch.bfloat16: 1.5, torch.float: 1.2}
127
+ # test matrix has inputs centered around 0 and lower precision from
128
+ # bfloat16 tends to accumulate and can greatly inflate rtol
129
+ # since outputs are also very close to 0
130
+ rtols = {torch.half: 1e-1, torch.bfloat16: 1e4, torch.float: 2e1}
131
+ torch.testing.assert_close(output,
132
+ ref_output,
133
+ atol=atols[dtype],
134
+ rtol=rtols[dtype])
vllm/_custom_ops.py
+2
-2
@@ -436,7 +436,7 @@ if hasattr(torch.ops._C, "ggml_dequantize"):
436
quant_type: int,
437
row: torch.SymInt,
438
) -> torch.Tensor:
439
- return torch.empty((1, row), dtype=torch.float16, device=W.device)
439
+ return torch.empty((1, row), dtype=X.dtype, device=W.device)
440
441
@register_fake("_C::ggml_mul_mat_a8")
442
def _ggml_mul_mat_a8_fake(
@@ -446,7 +446,7 @@ if hasattr(torch.ops._C, "ggml_dequantize"):
446
row: torch.SymInt,
447
) -> torch.Tensor:
448
batch = X.size(0)
449
- return torch.empty((batch, row), dtype=torch.float16, device=W.device)
449
+ return torch.empty((batch, row), dtype=X.dtype, device=W.device)
450
451
452
# cutlass
vllm/model_executor/layers/quantization/gguf.py
+3
-2
@@ -32,7 +32,7 @@ class GGUFConfig(QuantizationConfig):
32
return "gguf"
33
34
def get_supported_act_dtypes(self) -> List[torch.dtype]:
35
- return [torch.half]
35
+ return [torch.half, torch.bfloat16, torch.float32]
36
37
@classmethod
38
def get_min_capability(cls) -> int:
@@ -134,6 +134,7 @@ class GGUFLinearMethod(LinearMethodBase):
134
output_partition_sizes: List[int], input_size: int,
135
output_size: int, params_dtype: torch.dtype,
136
**extra_weight_attrs):
137
+ self.params_dtype = params_dtype
138
output_size_per_partition = sum(output_partition_sizes)
139
140
tensor_shape = (output_size_per_partition, input_size_per_partition)
@@ -326,7 +327,7 @@ class GGUFEmbeddingMethod(GGUFLinearMethod):
327
x_flat = x.flatten()
328
quant = torch.index_select(qweight, dim=0, index=x_flat)
329
dequant = ops.ggml_dequantize(quant, qweight_type, hidden_size,
329
- x_flat.shape[0])
330
+ x_flat.shape[0]).to(self.params_dtype)
331
return dequant.view(*x.shape, hidden_size)
332
333