Skip to content

Commit 873c825

Browse files
qnixsynapseqnixsynapse
authored
sycl: disable Q1_0 in backend and cleanup unused variables (ggml-org#21807)
1 parent 82764d8 commit 873c825

6 files changed

Lines changed: 18 additions & 12 deletions

File tree

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -488,7 +488,7 @@ static void dequantize_row_nvfp4_sycl(const void * vx, dst_t * y, const int64_t
488488
const int nb = k / QK_NVFP4;
489489
stream->parallel_for(
490490
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)),
491-
[=](sycl::nd_item<3> item_ct1) {
491+
[=](sycl::nd_item<3> /*item_ct1*/) {
492492
dequantize_block_nvfp4(vx, y, k);
493493
});
494494
}

‎ggml/src/ggml-sycl/dequantize.hpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -14,6 +14,7 @@
1414
#define GGML_SYCL_DEQUANTIZE_HPP
1515

1616
#include "common.hpp"
17+
#include "convert.hpp"
1718

1819
typedef void (*dequantize_kernel_t)(const void * vx, const int64_t ib, const int iqs, dfloat2 & v);
1920
typedef void (*dequantize_kernel_t_reorder)(const void *d, const int64_t ib, const void *qs,

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -355,7 +355,7 @@ static void acc_f32_sycl(const float *x, const float *y, float *dst,
355355
const int num_blocks = (n_elements + SYCL_ACC_BLOCK_SIZE - 1) / SYCL_ACC_BLOCK_SIZE;
356356
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE),
357357
sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE)),
358-
[=](sycl::nd_item<3> item_ct1) {
358+
[=](sycl::nd_item<3> /*item_ct1*/) {
359359
acc_f32(x, y, dst, n_elements, ne10, ne11, ne12, ne13, s1, s2, s3, offset);
360360
});
361361
}

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

Lines changed: 4 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -176,14 +176,12 @@ static void launch_gated_delta_net(const float * q_d,
176176
const sycl::uint3 neqk1_magic = init_fastdiv_values(neqk1);
177177
const sycl::uint3 rq3_magic = init_fastdiv_values(rq3);
178178

179-
int cc = ggml_sycl_info().devices[ggml_sycl_get_device()].cc;
180-
181179
switch (S_v) {
182180
case 16:
183181
{
184182
constexpr int sv = 16;
185183
stream->parallel_for(sycl::nd_range<3>(grid_dims * block_dims, block_dims),
186-
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
184+
[=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
187185
gated_delta_net_sycl<sv, KDA>(q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens,
188186
n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2,
189187
sb3, neqk1_magic, rq3_magic, scale);
@@ -194,7 +192,7 @@ static void launch_gated_delta_net(const float * q_d,
194192
{
195193
constexpr int sv = 32;
196194
stream->parallel_for(sycl::nd_range<3>(grid_dims * block_dims, block_dims),
197-
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
195+
[=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
198196
gated_delta_net_sycl<sv, KDA>(q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens,
199197
n_seqs, sq1, sq2, sq3, sv1, sv2, sv3, sb1, sb2,
200198
sb3, neqk1_magic, rq3_magic, scale);
@@ -205,7 +203,7 @@ static void launch_gated_delta_net(const float * q_d,
205203
{
206204
constexpr int sv = 64;
207205
stream->parallel_for(sycl::nd_range<3>(grid_dims * block_dims, block_dims),
208-
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
206+
[=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
209207
gated_delta_net_sycl<sv, KDA>(
210208
q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2,
211209
sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale);
@@ -217,7 +215,7 @@ static void launch_gated_delta_net(const float * q_d,
217215
{
218216
constexpr int sv = 128;
219217
stream->parallel_for(sycl::nd_range<3>(grid_dims * block_dims, block_dims),
220-
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
218+
[=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
221219
gated_delta_net_sycl<sv, KDA>(
222220
q_d, k_d, v_d, g_d, b_d, s_d, dst_d, H, n_tokens, n_seqs, sq1, sq2,
223221
sq3, sv1, sv2, sv3, sb1, sb2, sb3, neqk1_magic, rq3_magic, scale);

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

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4727,12 +4727,19 @@ static bool ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, const g
47274727
struct ggml_tensor * a = op->src[0];
47284728
struct ggml_tensor * b = op->src[1];
47294729

4730+
// disable Q1_0 until implementation
4731+
if (a->type == GGML_TYPE_Q1_0 || b->type == GGML_TYPE_Q1_0) {
4732+
return false;
4733+
}
4734+
47304735
if (a->ne[3] != b->ne[3]) {
47314736
return false;
47324737
}
47334738

47344739
ggml_type src0_type = op->src[0]->type;
47354740

4741+
4742+
47364743
// TODO: The configuration below needs more work to be supported with oneDNN
47374744
if (ggml_is_permuted(a) && !ggml_is_contiguous(a) &&
47384745
a->ne[2] > 1 && a->ne[3] > 1 && src0_type == GGML_TYPE_F16) {

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -272,7 +272,7 @@ static void upscale_f32_sycl(const float * x,
272272
sycl::nd_range<3>(
273273
sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_UPSCALE_BLOCK_SIZE),
274274
sycl::range<3>(1, 1, SYCL_UPSCALE_BLOCK_SIZE)),
275-
[=](sycl::nd_item<3> item_ct1) {
275+
[=](sycl::nd_item<3> /*item_ct1*/) {
276276
upscale_f32(x, dst, nb00, nb01, nb02, nb03, ne10, ne11, ne12, ne13, sf0, sf1, sf2, sf3);
277277
});
278278
}
@@ -304,7 +304,7 @@ static void upscale_f32_bilinear_sycl(const float * x,
304304
sycl::nd_range<3>(
305305
sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_UPSCALE_BLOCK_SIZE),
306306
sycl::range<3>(1, 1, SYCL_UPSCALE_BLOCK_SIZE)),
307-
[=](sycl::nd_item<3> item_ct1) {
307+
[=](sycl::nd_item<3> /*item_ct1*/) {
308308
upscale_f32_bilinear_antialias(
309309
x, dst, nb00, nb01, nb02, nb03, ne00_src, ne01_src, ne10_dst, ne11_dst,
310310
ne12_dst, ne13_dst, sf0, sf1, sf2, sf3, pixel_offset);
@@ -314,7 +314,7 @@ static void upscale_f32_bilinear_sycl(const float * x,
314314
sycl::nd_range<3>(
315315
sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_UPSCALE_BLOCK_SIZE),
316316
sycl::range<3>(1, 1, SYCL_UPSCALE_BLOCK_SIZE)),
317-
[=](sycl::nd_item<3> item_ct1) {
317+
[=](sycl::nd_item<3> /*item_ct1*/) {
318318
upscale_f32_bilinear(
319319
x, dst, nb00, nb01, nb02, nb03, ne00_src, ne01_src, ne10_dst, ne11_dst, ne12_dst,
320320
ne13_dst, sf0, sf1, sf2, sf3, pixel_offset);
@@ -349,7 +349,7 @@ static void upscale_f32_bicubic_sycl(const float * x,
349349
sycl::nd_range<3>(
350350
sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_UPSCALE_BLOCK_SIZE),
351351
sycl::range<3>(1, 1, SYCL_UPSCALE_BLOCK_SIZE)),
352-
[=](sycl::nd_item<3> item_ct1) {
352+
[=](sycl::nd_item<3> /*item_ct1*/) {
353353
upscale_f32_bicubic(
354354
x, dst, nb00, nb01, nb02, nb03, ne00_src, ne01_src, ne10_dst, ne11_dst,
355355
ne12_dst, ne13_dst, sf0, sf1, sf2, sf3, pixel_offset);

0 commit comments

Comments
 (0)