Skip to content

Commit 83ed24a

Browse files
qnixsynapseabhilash1910
qnixsynapse
and
abhilash1910
authored
SYCL: Reduce most of the compiler warnings (ggml-org#10748)
* Try to reduce some unused and typecast warnings * Reduce compiler warnings step 2 * add a newline at the end of the file * Initialize nreduce as size_t * [SYCL] Remove pragma directives from mmq.cpp * SYCL: mmq add condition to prevent blocks_per_tile_x_row variable from becoming 0 * SYCL softmax: Initialize nreduce as size_t * ggml-sycl.cpp: fix some trailing whitespaces * SYCL: remove the unused variables instead of commenting it out * SYCL poo2d kernel: set NAN for invalid pooling op * SYCL gemm.hpp: remove pragma directives * SYCL gemm.hpp: use const cast to properly support dnnl::memory * SYCL: wkv6 remove a comment * SYCL: clean comments step 2 * SYCL: clean comments and variables step 3 * SYCL: Use GGML_UNUSED for unused variables * SYCL: remove extra empty lines and a comment * Remove TODO * cleanup spaces * add a stdout for unsupported op * use sycl printf over fprintf * remove prints for CI * SYCL ggml-sycl: pool2D use sycl::nan and remove if-else block --------- Co-authored-by: Abhilash Majumder <[email protected]>
1 parent d583cd0 commit 83ed24a

File tree

17 files changed

+205
-187
lines changed

17 files changed

+205
-187
lines changed

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

Lines changed: 2 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,7 @@
1111
//
1212

1313
#include "common.hpp"
14+
#include "ggml-impl.h"
1415

1516
int get_current_device_id() {
1617
return dpct::dev_mgr::instance().current_device_id();
@@ -28,11 +29,7 @@ void* ggml_sycl_host_malloc(size_t size) try {
2829

2930
if (err != 0) {
3031
// clear the error
31-
fprintf(
32-
stderr,
33-
"WARNING: failed to allocate %.2f MB of pinned memory: %s\n",
34-
size / 1024.0 / 1024.0,
35-
"syclGetErrorString is not supported");
32+
GGML_LOG_ERROR("WARNING: failed to allocate %.2f MB of pinned memory: %s\n", size / 1024.0 / 1024.0, "syclGetErrorString is not supported");
3633
return nullptr;
3734
}
3835

@@ -66,18 +63,12 @@ int64_t downsample_sycl_global_range(int64_t accumulate_block_num, int64_t block
6663
void ggml_sycl_op_flatten(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
6764
const ggml_tensor *src1, ggml_tensor *dst,
6865
const ggml_sycl_op_flatten_t op) try {
69-
const int64_t nrows0 = ggml_nrows(src0);
7066

7167
const bool use_src1 = src1 != nullptr;
72-
const int64_t nrows1 = use_src1 ? ggml_nrows(src1) : 1;
7368

7469
GGML_ASSERT(!use_src1 || src1->backend != GGML_BACKEND_TYPE_GPU_SPLIT);
7570
GGML_ASSERT( dst->backend != GGML_BACKEND_TYPE_GPU_SPLIT);
7671

77-
ggml_tensor_extra_gpu * src0_extra = (ggml_tensor_extra_gpu *) src0->extra;
78-
ggml_tensor_extra_gpu * src1_extra = use_src1 ? (ggml_tensor_extra_gpu *) src1->extra : nullptr;
79-
ggml_tensor_extra_gpu * dst_extra = (ggml_tensor_extra_gpu *) dst->extra;
80-
8172
// dd = data device
8273
float * src0_ddf = (float *) src0->data;
8374
float * src1_ddf = use_src1 ? (float *) src1->data : nullptr;

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -626,6 +626,7 @@ struct bin_bcast_sycl {
626626
});
627627
}
628628
}
629+
GGML_UNUSED(ctx);
629630
}
630631
};
631632

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -47,7 +47,7 @@ static void concat_f32_dim1(const float *x, const float *y, float *dst,
4747
// operation
4848
int offset_dst = nidx + item_ct1.get_group(1) * ne0 +
4949
item_ct1.get_group(0) * ne0 * item_ct1.get_group_range(1);
50-
if (item_ct1.get_group(1) < ne01) { // src0
50+
if (item_ct1.get_group(1) < (size_t) ne01) { // src0
5151
int offset_src =
5252
nidx + item_ct1.get_group(1) * ne0 + item_ct1.get_group(0) * ne0 * ne01;
5353
dst[offset_dst] = x[offset_src];
@@ -70,7 +70,7 @@ static void concat_f32_dim2(const float *x, const float *y, float *dst,
7070
// operation
7171
int offset_dst = nidx + item_ct1.get_group(1) * ne0 +
7272
item_ct1.get_group(0) * ne0 * item_ct1.get_group_range(1);
73-
if (item_ct1.get_group(0) < ne02) { // src0
73+
if (item_ct1.get_group(0) < (size_t) ne02) { // src0
7474
int offset_src = nidx + item_ct1.get_group(1) * ne0 +
7575
item_ct1.get_group(0) * ne0 * item_ct1.get_group_range(1);
7676
dst[offset_dst] = x[offset_src];

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -424,7 +424,7 @@ static void convert_unary(const void * __restrict__ vx, dst_t * __restrict__ y,
424424
const int64_t global_id = item_ct1.get_local_id(2) + work_group_size * item_ct1.get_group(2);
425425

426426
// make each work-item deal with more elements since sycl global range can not exceed max int
427-
const src_t * x = (src_t *) vx;
427+
const src_t * x = (const src_t *) vx;
428428
for (int64_t i = global_id; i < k; i += work_group_size * item_ct1.get_group_range(2)) {
429429
y[i] = x[i];
430430
}

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

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1015,9 +1015,9 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
10151015
break;
10161016
}
10171017

1018-
(void) src1;
1019-
(void) dst;
1020-
(void) src1_ddq_i;
1021-
(void) src1_ncols;
1022-
(void) src1_padded_row_size;
1018+
GGML_UNUSED(src1);
1019+
GGML_UNUSED(dst);
1020+
GGML_UNUSED(src1_ddq_i);
1021+
GGML_UNUSED(src1_ncols);
1022+
GGML_UNUSED(src1_padded_row_size);
10231023
}

‎ggml/src/ggml-sycl/dpct/helper.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1237,7 +1237,7 @@ namespace dpct
12371237
12381238
std::map<byte_t *, allocation>::iterator get_map_iterator(const void *ptr)
12391239
{
1240-
auto it = m_map.upper_bound((byte_t *)ptr);
1240+
auto it = m_map.upper_bound(const_cast<byte_t *>(reinterpret_cast<const byte_t *>(ptr)));
12411241
if (it == m_map.end())
12421242
{
12431243
// Not a virtual pointer.

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

Lines changed: 80 additions & 61 deletions
Original file line numberDiff line numberDiff line change
@@ -237,7 +237,7 @@ void upscale_f32(const float *x, float *dst, const int nb00, const int nb01,
237237
int i02 = i12 / sf2;
238238
int i03 = i13 / sf3;
239239

240-
dst[index] = *(float *)((char *)x + i03 * nb03 + i02 * nb02 + i01 * nb01 + i00 * nb00);
240+
dst[index] = *(const float *)((const char *)x + i03 * nb03 + i02 * nb02 + i01 * nb01 + i00 * nb00);
241241
}
242242

243243
void pad_f32(const float *x, float *dst, const int ne0, const int ne00, const int ne01, const int ne02,
@@ -251,8 +251,7 @@ void pad_f32(const float *x, float *dst, const int ne0, const int ne00, const i
251251
// operation
252252
int offset_dst = nidx + item_ct1.get_group(1) * ne0 +
253253
item_ct1.get_group(0) * ne0 * item_ct1.get_group_range(1);
254-
if (nidx < ne00 && item_ct1.get_group(1) < ne01 &&
255-
item_ct1.get_group(0) < ne02) {
254+
if (nidx < ne00 && item_ct1.get_group(1) < (size_t) ne01 && item_ct1.get_group(0) < (size_t) ne02) {
256255
int offset_src = nidx + item_ct1.get_group(1) * ne00 +
257256
item_ct1.get_group(0) * ne00 * ne01;
258257
dst[offset_dst] = x[offset_src];
@@ -520,9 +519,10 @@ inline void ggml_sycl_op_silu(ggml_backend_sycl_context & ctx, const ggml_tensor
520519

521520
silu_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
522521

523-
(void) src1;
524-
(void) dst;
525-
(void) src1_dd;
522+
GGML_UNUSED(src1);
523+
GGML_UNUSED(dst);
524+
GGML_UNUSED(src1_dd);
525+
GGML_UNUSED(ctx);
526526
}
527527

528528
inline void ggml_sycl_op_gelu(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
@@ -535,9 +535,10 @@ inline void ggml_sycl_op_gelu(ggml_backend_sycl_context & ctx, const ggml_tensor
535535

536536
gelu_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
537537

538-
(void) src1;
539-
(void) dst;
540-
(void) src1_dd;
538+
GGML_UNUSED(src1);
539+
GGML_UNUSED(dst);
540+
GGML_UNUSED(src1_dd);
541+
GGML_UNUSED(ctx);
541542
}
542543
inline void ggml_sycl_op_gelu_quick(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
543544
const ggml_tensor *src1, ggml_tensor *dst,
@@ -550,9 +551,10 @@ inline void ggml_sycl_op_gelu_quick(ggml_backend_sycl_context & ctx, const ggml_
550551

551552
gelu_quick_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
552553

553-
(void) src1;
554-
(void) dst;
555-
(void) src1_dd;
554+
GGML_UNUSED(src1);
555+
GGML_UNUSED(dst);
556+
GGML_UNUSED(src1_dd);
557+
GGML_UNUSED(ctx);
556558
}
557559

558560
inline void ggml_sycl_op_tanh(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
@@ -564,9 +566,10 @@ inline void ggml_sycl_op_tanh(ggml_backend_sycl_context & ctx, const ggml_tensor
564566
GGML_ASSERT( dst->type == GGML_TYPE_F32);
565567
tanh_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
566568

567-
(void) src1;
568-
(void) dst;
569-
(void) src1_dd;
569+
GGML_UNUSED(src1);
570+
GGML_UNUSED(dst);
571+
GGML_UNUSED(src1_dd);
572+
GGML_UNUSED(ctx);
570573
}
571574

572575
inline void ggml_sycl_op_relu(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
@@ -579,9 +582,10 @@ inline void ggml_sycl_op_relu(ggml_backend_sycl_context & ctx, const ggml_tensor
579582

580583
relu_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
581584

582-
(void) src1;
583-
(void) dst;
584-
(void) src1_dd;
585+
GGML_UNUSED(src1);
586+
GGML_UNUSED(dst);
587+
GGML_UNUSED(src1_dd);
588+
GGML_UNUSED(ctx);
585589
}
586590

587591
inline void ggml_sycl_op_hardsigmoid(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -595,9 +599,10 @@ inline void ggml_sycl_op_hardsigmoid(ggml_backend_sycl_context & ctx, const ggml
595599

596600
hardsigmoid_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
597601

598-
(void) src1;
599-
(void) dst;
600-
(void) src1_dd;
602+
GGML_UNUSED(src1);
603+
GGML_UNUSED(dst);
604+
GGML_UNUSED(src1_dd);
605+
GGML_UNUSED(ctx);
601606
}
602607

603608
inline void ggml_sycl_op_hardswish(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -610,9 +615,10 @@ inline void ggml_sycl_op_hardswish(ggml_backend_sycl_context & ctx, const ggml_t
610615

611616
hardswish_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
612617

613-
(void) src1;
614-
(void) dst;
615-
(void) src1_dd;
618+
GGML_UNUSED(src1);
619+
GGML_UNUSED(dst);
620+
GGML_UNUSED(src1_dd);
621+
GGML_UNUSED(ctx);
616622
}
617623

618624
inline void ggml_sycl_op_exp(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -625,9 +631,10 @@ inline void ggml_sycl_op_exp(ggml_backend_sycl_context & ctx, const ggml_tensor
625631

626632
exp_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
627633

628-
(void) src1;
629-
(void) dst;
630-
(void) src1_dd;
634+
GGML_UNUSED(src1);
635+
GGML_UNUSED(dst);
636+
GGML_UNUSED(src1_dd);
637+
GGML_UNUSED(ctx);
631638
}
632639

633640
inline void ggml_sycl_op_log(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -640,9 +647,10 @@ inline void ggml_sycl_op_log(ggml_backend_sycl_context & ctx, const ggml_tensor
640647

641648
log_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
642649

643-
(void) src1;
644-
(void) dst;
645-
(void) src1_dd;
650+
GGML_UNUSED(src1);
651+
GGML_UNUSED(dst);
652+
GGML_UNUSED(src1_dd);
653+
GGML_UNUSED(ctx);
646654
}
647655

648656
inline void ggml_sycl_op_sigmoid(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -655,9 +663,10 @@ inline void ggml_sycl_op_sigmoid(ggml_backend_sycl_context & ctx, const ggml_ten
655663

656664
sigmoid_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
657665

658-
(void) src1;
659-
(void) dst;
660-
(void) src1_dd;
666+
GGML_UNUSED(src1);
667+
GGML_UNUSED(dst);
668+
GGML_UNUSED(src1_dd);
669+
GGML_UNUSED(ctx);
661670
}
662671

663672
inline void ggml_sycl_op_sqrt(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -670,9 +679,10 @@ inline void ggml_sycl_op_sqrt(ggml_backend_sycl_context & ctx, const ggml_tensor
670679

671680
sqrt_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
672681

673-
(void) src1;
674-
(void) dst;
675-
(void) src1_dd;
682+
GGML_UNUSED(src1);
683+
GGML_UNUSED(dst);
684+
GGML_UNUSED(src1_dd);
685+
GGML_UNUSED(ctx);
676686
}
677687

678688
inline void ggml_sycl_op_sin(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -685,9 +695,10 @@ inline void ggml_sycl_op_sin(ggml_backend_sycl_context & ctx, const ggml_tensor
685695

686696
sin_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
687697

688-
(void) src1;
689-
(void) dst;
690-
(void) src1_dd;
698+
GGML_UNUSED(src1);
699+
GGML_UNUSED(dst);
700+
GGML_UNUSED(src1_dd);
701+
GGML_UNUSED(ctx);
691702
}
692703

693704
inline void ggml_sycl_op_cos(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -700,9 +711,10 @@ inline void ggml_sycl_op_cos(ggml_backend_sycl_context & ctx, const ggml_tensor
700711

701712
cos_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
702713

703-
(void) src1;
704-
(void) dst;
705-
(void) src1_dd;
714+
GGML_UNUSED(src1);
715+
GGML_UNUSED(dst);
716+
GGML_UNUSED(src1_dd);
717+
GGML_UNUSED(ctx);
706718
}
707719

708720
inline void ggml_sycl_op_step(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -715,9 +727,10 @@ inline void ggml_sycl_op_step(ggml_backend_sycl_context & ctx, const ggml_tensor
715727

716728
step_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
717729

718-
(void) src1;
719-
(void) dst;
720-
(void) src1_dd;
730+
GGML_UNUSED(src1);
731+
GGML_UNUSED(dst);
732+
GGML_UNUSED(src1_dd);
733+
GGML_UNUSED(ctx);
721734
}
722735

723736
inline void ggml_sycl_op_neg(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -730,9 +743,10 @@ inline void ggml_sycl_op_neg(ggml_backend_sycl_context & ctx, const ggml_tensor
730743

731744
neg_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
732745

733-
(void) src1;
734-
(void) dst;
735-
(void) src1_dd;
746+
GGML_UNUSED(src1);
747+
GGML_UNUSED(dst);
748+
GGML_UNUSED(src1_dd);
749+
GGML_UNUSED(ctx);
736750
}
737751

738752
inline void ggml_sycl_op_leaky_relu(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -749,9 +763,10 @@ inline void ggml_sycl_op_leaky_relu(ggml_backend_sycl_context & ctx, const ggml_
749763

750764
leaky_relu_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), negative_slope, main_stream);
751765

752-
(void) src1;
753-
(void) dst;
754-
(void) src1_dd;
766+
GGML_UNUSED(src1);
767+
GGML_UNUSED(dst);
768+
GGML_UNUSED(src1_dd);
769+
GGML_UNUSED(ctx);
755770
}
756771

757772
inline void ggml_sycl_op_sqr(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
@@ -764,9 +779,10 @@ inline void ggml_sycl_op_sqr(ggml_backend_sycl_context & ctx, const ggml_tensor
764779

765780
sqr_f32_sycl(src0_dd, dst_dd, ggml_nelements(src0), main_stream);
766781

767-
(void) src1;
768-
(void) dst;
769-
(void) src1_dd;
782+
GGML_UNUSED(src1);
783+
GGML_UNUSED(dst);
784+
GGML_UNUSED(src1_dd);
785+
GGML_UNUSED(ctx);
770786
}
771787

772788
inline void ggml_sycl_op_upscale(ggml_backend_sycl_context & ctx, const ggml_tensor *src0,
@@ -787,9 +803,10 @@ inline void ggml_sycl_op_upscale(ggml_backend_sycl_context & ctx, const ggml_ten
787803
dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3], sf0, sf1, sf2, sf3,
788804
main_stream);
789805

790-
(void) src1;
791-
(void) dst;
792-
(void) src1_dd;
806+
GGML_UNUSED(src1);
807+
GGML_UNUSED(dst);
808+
GGML_UNUSED(src1_dd);
809+
GGML_UNUSED(ctx);
793810
}
794811

795812
inline void ggml_sycl_op_pad(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
@@ -805,9 +822,10 @@ inline void ggml_sycl_op_pad(ggml_backend_sycl_context & ctx, const ggml_tensor
805822
src0->ne[0], src0->ne[1], src0->ne[2],
806823
dst->ne[0], dst->ne[1], dst->ne[2], main_stream);
807824

808-
(void) src1;
809-
(void) dst;
810-
(void) src1_dd;
825+
GGML_UNUSED(src1);
826+
GGML_UNUSED(dst);
827+
GGML_UNUSED(src1_dd);
828+
GGML_UNUSED(ctx);
811829
}
812830

813831
inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
@@ -827,7 +845,8 @@ inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, const ggml_tensor
827845

828846
acc_f32_sycl(src0_dd, src1_dd, dst_dd, ggml_nelements(dst), src1->ne[0], src1->ne[1], src1->ne[2], nb1, nb2, offset, main_stream);
829847

830-
(void) dst;
848+
GGML_UNUSED(dst);
849+
GGML_UNUSED(ctx);
831850
}
832851

833852
inline void ggml_sycl_op_add(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,

0 commit comments

Comments
 (0)