Skip to content

Commit 7657ec3

Browse files
CISCCISC
authored
test quantized path
warnings--
1 parent 9a7bff3 commit 7657ec3

File tree

4 files changed

+21
-20
lines changed

4 files changed

+21
-20
lines changed

‎ggml/src/ggml-cuda/set-rows.cu‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -181,7 +181,7 @@ void ggml_cuda_op_set_rows(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
181181
nb1, nb2, nb3,
182182
stream
183183
);
184-
} else if (src1->type == GGML_TYPE_I32) {
184+
} else {
185185
set_rows_cuda(
186186
src0_d, (const int32_t *)src1->data, (float*)dst->data,
187187
ne00, ne01, ne02, ne03,
@@ -203,7 +203,7 @@ void ggml_cuda_op_set_rows(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
203203
nb1, nb2, nb3,
204204
stream
205205
);
206-
} else if (src1->type == GGML_TYPE_I32) {
206+
} else {
207207
set_rows_cuda(
208208
src0_d, (const int32_t *)src1->data, (half*)dst->data,
209209
ne00, ne01, ne02, ne03,
@@ -225,7 +225,7 @@ void ggml_cuda_op_set_rows(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
225225
nb1, nb2, nb3,
226226
stream
227227
);
228-
} else if (src1->type == GGML_TYPE_I32) {
228+
} else {
229229
set_rows_cuda(
230230
src0_d, (const int32_t *)src1->data, (nv_bfloat16*)dst->data,
231231
ne00, ne01, ne02, ne03,
@@ -247,7 +247,7 @@ void ggml_cuda_op_set_rows(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
247247
nb1, nb2, nb3,
248248
stream
249249
);
250-
} else if (src1->type == GGML_TYPE_I32) {
250+
} else {
251251
set_rows_cuda_quant<int32_t, block_q4_0, QK4_0, quantize_f32_q4_0_block>(
252252
src0_d, (const int32_t *)src1->data, (block_q4_0*)dst->data,
253253
ne00, ne01, ne02, ne03,
@@ -269,7 +269,7 @@ void ggml_cuda_op_set_rows(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
269269
nb1, nb2, nb3,
270270
stream
271271
);
272-
} else if (src1->type == GGML_TYPE_I32) {
272+
} else {
273273
set_rows_cuda_quant<int32_t, block_q4_1, QK4_1, quantize_f32_q4_1_block>(
274274
src0_d, (const int32_t *)src1->data, (block_q4_1*)dst->data,
275275
ne00, ne01, ne02, ne03,
@@ -291,7 +291,7 @@ void ggml_cuda_op_set_rows(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
291291
nb1, nb2, nb3,
292292
stream
293293
);
294-
} else if (src1->type == GGML_TYPE_I32) {
294+
} else {
295295
set_rows_cuda_quant<int32_t, block_q5_0, QK5_0, quantize_f32_q5_0_block>(
296296
src0_d, (const int32_t *)src1->data, (block_q5_0*)dst->data,
297297
ne00, ne01, ne02, ne03,
@@ -313,7 +313,7 @@ void ggml_cuda_op_set_rows(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
313313
nb1, nb2, nb3,
314314
stream
315315
);
316-
} else if (src1->type == GGML_TYPE_I32) {
316+
} else {
317317
set_rows_cuda_quant<int32_t, block_q5_1, QK5_1, quantize_f32_q5_1_block>(
318318
src0_d, (const int32_t *)src1->data, (block_q5_1*)dst->data,
319319
ne00, ne01, ne02, ne03,
@@ -335,7 +335,7 @@ void ggml_cuda_op_set_rows(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
335335
nb1, nb2, nb3,
336336
stream
337337
);
338-
} else if (src1->type == GGML_TYPE_I32) {
338+
} else {
339339
set_rows_cuda_quant<int32_t, block_q8_0, QK8_0, quantize_f32_q8_0_block>(
340340
src0_d, (const int32_t *)src1->data, (block_q8_0*)dst->data,
341341
ne00, ne01, ne02, ne03,
@@ -357,7 +357,7 @@ void ggml_cuda_op_set_rows(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
357357
nb1, nb2, nb3,
358358
stream
359359
);
360-
} else if (src1->type == GGML_TYPE_I32) {
360+
} else {
361361
set_rows_cuda_quant<int32_t, block_iq4_nl, QK4_NL, quantize_f32_iq4_nl_block>(
362362
src0_d, (const int32_t *)src1->data, (block_iq4_nl*)dst->data,
363363
ne00, ne01, ne02, ne03,

‎ggml/src/ggml-opencl/ggml-opencl.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4111,14 +4111,14 @@ static void ggml_cl_set_rows(ggml_backend_t backend, const ggml_tensor * src0, c
41114111
case GGML_TYPE_F32:
41124112
if (src1->type == GGML_TYPE_I64) {
41134113
kernel = backend_ctx->kernel_set_rows_f32_i64;
4114-
} else if (src1->type == GGML_TYPE_I32) {
4114+
} else {
41154115
kernel = backend_ctx->kernel_set_rows_f32_i32;
41164116
}
41174117
break;
41184118
case GGML_TYPE_F16:
41194119
if (src1->type == GGML_TYPE_I64) {
41204120
kernel = backend_ctx->kernel_set_rows_f16_i64;
4121-
} else if (src1->type == GGML_TYPE_I32) {
4121+
} else {
41224122
kernel = backend_ctx->kernel_set_rows_f16_i32;
41234123
}
41244124
break;

‎ggml/src/ggml-sycl/set_rows.cpp‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -171,7 +171,7 @@ void ggml_sycl_op_set_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
171171
sizeof(float), sizeof(float),
172172
stream
173173
);
174-
} else if (src1->type == GGML_TYPE_I32) {
174+
} else {
175175
set_rows_sycl<float, int32_t, float>(
176176
(const char *)src0->data, (const int32_t *)src1->data, (char *)dst->data,
177177
ne00, ne01, ne02, ne03,
@@ -197,7 +197,7 @@ void ggml_sycl_op_set_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
197197
sizeof(float), sizeof(sycl::half),
198198
stream
199199
);
200-
} else if (src1->type == GGML_TYPE_I32) {
200+
} else {
201201
set_rows_sycl<float, int32_t, sycl::half>(
202202
(const char *)src0->data, (const int32_t *)src1->data, (char *)dst->data,
203203
ne00, ne01, ne02, ne03,
@@ -222,7 +222,7 @@ void ggml_sycl_op_set_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
222222
sizeof(float), sizeof(sycl::ext::oneapi::bfloat16),
223223
stream
224224
);
225-
} else if (src1->type == GGML_TYPE_I32) {
225+
} else {
226226
set_rows_sycl<float, int32_t, sycl::ext::oneapi::bfloat16>(
227227
(const char *)src0->data, (const int32_t *)src1->data, (char *)dst->data,
228228
ne00, ne01, ne02, ne03,
@@ -238,42 +238,42 @@ void ggml_sycl_op_set_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
238238
case GGML_TYPE_Q8_0:
239239
if (src1->type == GGML_TYPE_I64) {
240240
set_rows_sycl_q<int64_t, block_q8_0, QK8_0, cpy_blck_f32_q8_0>((const char *)src0->data, (const int64_t *)src1->data, (block_q8_0 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
241-
} else if (src1->type == GGML_TYPE_I32) {
241+
} else {
242242
set_rows_sycl_q<int32_t, block_q8_0, QK8_0, cpy_blck_f32_q8_0>((const char *)src0->data, (const int32_t *)src1->data, (block_q8_0 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
243243
}
244244
break;
245245
case GGML_TYPE_Q5_1:
246246
if (src1->type == GGML_TYPE_I64) {
247247
set_rows_sycl_q<int64_t, block_q5_1, QK5_1, cpy_blck_f32_q5_1>((const char *)src0->data, (const int64_t *)src1->data, (block_q5_1 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
248-
} else if (src1->type == GGML_TYPE_I32) {
248+
} else {
249249
set_rows_sycl_q<int32_t, block_q5_1, QK5_1, cpy_blck_f32_q5_1>((const char *)src0->data, (const int32_t *)src1->data, (block_q5_1 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
250250
}
251251
break;
252252
case GGML_TYPE_Q5_0:
253253
if (src1->type == GGML_TYPE_I64) {
254254
set_rows_sycl_q<int64_t, block_q5_0, QK5_0, cpy_blck_f32_q5_0>((const char *)src0->data, (const int64_t *)src1->data, (block_q5_0 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
255-
} else if (src1->type == GGML_TYPE_I32) {
255+
} else {
256256
set_rows_sycl_q<int32_t, block_q5_0, QK5_0, cpy_blck_f32_q5_0>((const char *)src0->data, (const int32_t *)src1->data, (block_q5_0 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
257257
}
258258
break;
259259
case GGML_TYPE_Q4_1:
260260
if (src1->type == GGML_TYPE_I64) {
261261
set_rows_sycl_q<int64_t, block_q4_1, QK4_1, cpy_blck_f32_q4_1>((const char *)src0->data, (const int64_t *)src1->data, (block_q4_1 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
262-
} else if (src1->type == GGML_TYPE_I32) {
262+
} else {
263263
set_rows_sycl_q<int32_t, block_q4_1, QK4_1, cpy_blck_f32_q4_1>((const char *)src0->data, (const int32_t *)src1->data, (block_q4_1 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
264264
}
265265
break;
266266
case GGML_TYPE_Q4_0:
267267
if (src1->type == GGML_TYPE_I64) {
268268
set_rows_sycl_q<int64_t, block_q4_0, QK4_0, cpy_blck_f32_q4_0>((const char *)src0->data, (const int64_t *)src1->data, (block_q4_0 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
269-
} else if (src1->type == GGML_TYPE_I32) {
269+
} else {
270270
set_rows_sycl_q<int32_t, block_q4_0, QK4_0, cpy_blck_f32_q4_0>((const char *)src0->data, (const int32_t *)src1->data, (block_q4_0 *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
271271
}
272272
break;
273273
case GGML_TYPE_IQ4_NL:
274274
if (src1->type == GGML_TYPE_I64) {
275275
set_rows_sycl_q<int64_t, block_iq4_nl, QK4_NL, cpy_blck_f32_iq4_nl>((const char *)src0->data, (const int64_t *)src1->data, (block_iq4_nl *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
276-
} else if (src1->type == GGML_TYPE_I32) {
276+
} else {
277277
set_rows_sycl_q<int32_t, block_iq4_nl, QK4_NL, cpy_blck_f32_iq4_nl>((const char *)src0->data, (const int32_t *)src1->data, (block_iq4_nl *)dst->data, ne00, ne01, ne02, ne03, ne10, ne11, ne12, ne13, nb00, nb01, nb02, nb03, nb10, nb11, nb12, nb13, nb1, nb2, nb3, stream);
278278
}
279279
break;

‎tests/test-backend-ops.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5671,6 +5671,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
56715671

56725672
test_cases.emplace_back(new test_set_rows(GGML_TYPE_F32, GGML_TYPE_I64, { 1, 8, 1, 3 }, { 1, 1 }, 2, false));
56735673
test_cases.emplace_back(new test_set_rows(GGML_TYPE_F32, GGML_TYPE_I32, { 1, 8, 1, 3 }, { 1, 1 }, 2, false));
5674+
test_cases.emplace_back(new test_set_rows(GGML_TYPE_Q8_0, GGML_TYPE_I32, { 256, 5, 1, 3 }, { 1, 1, }, 1, false));
56745675
for (ggml_type type : all_types) {
56755676
for (int b : {1, 7}) {
56765677
for (bool v : {false, true}) {

0 commit comments

Comments
 (0)