public_git
/
whisper.cpp
mirror of https://github.com/ggml-org/whisper.cpp.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

sycl: fix multiple warnings in compiling sycl backend (llama/26713) 

* Update norm.cpp

* Update helper.hpp

* Update im2col.cpp

* Update fattn-mkl.cpp

* Update element_wise.cpp

* Update fattn-mkl.cpp

* Update set_rows.cpp

* Update element_wise.cpp

* Update ggml-sycl.cpp

* Update ggml-sycl.cpp

* Update ggml-sycl.cpp

* Update ggml-sycl.cpp

* Update ggml-sycl.cpp

* Update norm.cpp

* Update CMakeLists.txt

* Update CMakeLists.txt

* Update CMakeLists.txt

* Update ggml-sycl.cpp
This commit is contained in:
Ian Faust 2026-08-21 07:15:40 +02:00 committed by Georgi Gerganov
parent 12137c316d
commit f19250e5ac
8 changed files with 28 additions and 34 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
ggml/src/ggml-cpu/CMakeLists.txt
View File

@ -737,8 +737,9 @@ function(ggml_add_cpu_backend_variant_impl tag_name)
set_target_properties(${GGML_CPU_NAME} PROPERTIES COMPILE_FLAGS "-msimd128")
endif()
if (CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM")
# The compiler automatically enables "-ffast-math" which can cause NaNs in tests due to "-fassociative-math"
target_compile_options(${GGML_CPU_NAME} PRIVATE "-fno-associative-math")
endif()
if (CMAKE_C_COMPILER_ID STREQUAL "IntelLLVM" OR CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM")
# The compiler automatically enables "-ffast-math" which can cause NaNs in tests due to "-fassociative-math"
target_compile_options(${GGML_CPU_NAME} PRIVATE "$<$<OR:$<COMPILE_LANG_AND_ID:C,IntelLLVM>,$<COMPILE_LANG_AND_ID:CXX,IntelLLVM>>:$<$<BOOL:${WIN32}>:/clang:>-fno-associative-math>")
endif()
endfunction()

2
ggml/src/ggml-sycl/dpct/helper.hpp
View File

@ -62,7 +62,7 @@
#define DPCT_UNUSED(x) (void)(x)
inline void _abort(const char * str) {
[[noreturn]] inline void _abort(const char * str) {
std::cerr << str << std::endl;
std::abort();
}

8
ggml/src/ggml-sycl/element_wise.cpp
View File

@ -10,7 +10,7 @@
(ITEM.get_local_range(IDX) * ITEM.get_group(IDX) + ITEM.get_local_id(IDX))
static void acc_f32(const char * x, const char * y, float * dst, const int64_t ne,
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3,
const int64_t ne0, const int64_t ne1, const int64_t ne2,
const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
@ -455,7 +455,7 @@ static void unary_mul_sycl(const T * x, const T * g, T * dst, const int64_t k, c
namespace ggml_sycl_detail {
static void acc_f32_sycl(const char *x, const char *y, float *dst,
const int64_t n_elements,
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3,
const int64_t ne0, const int64_t ne1, const int64_t ne2,
const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
const int64_t ne10, const int64_t ne11, const int64_t ne12, const int64_t ne13,
const int64_t nb10, const int64_t nb11, const int64_t nb12, const int64_t nb13,
@ -466,7 +466,7 @@ static void acc_f32_sycl(const char *x, const char *y, float *dst,
sycl::range<3>(1, 1, SYCL_ACC_BLOCK_SIZE)),
[=](sycl::nd_item<3> /*item_ct1*/) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
acc_f32(x, y, dst, n_elements,
ne0, ne1, ne2, ne3,
ne0, ne1, ne2,
nb00, nb01, nb02, nb03,
ne10, ne11, ne12, ne13,
nb10, nb11, nb12, nb13,
@ -970,7 +970,7 @@ static inline void ggml_sycl_op_acc(ggml_backend_sycl_context & ctx, ggml_tensor
const int64_t offset = (int64_t) ((const int32_t *) dst->op_params)[3] / (int64_t) sizeof(float);
ggml_sycl_detail::acc_f32_sycl(src0_d, src1_d, dst_d, ggml_nelements(dst),
dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3],
dst->ne[0], dst->ne[1], dst->ne[2],
src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
src1->nb[0], src1->nb[1], src1->nb[2], src1->nb[3],

9
ggml/src/ggml-sycl/fattn-mkl.cpp
View File

@ -43,7 +43,7 @@ static void mkl_fa_pack_q_fp16(
dpct::queue_ptr stream,
sycl::half * __restrict dst,
const float * __restrict q_src,
int n_queries, int n_query_rows, int DKQ,
int n_queries, int DKQ,
int gqa_ratio, int kvh_base_head,
float q_scale, int64_t q_row_stride, int64_t q_head_stride,
int64_t wg_size) {
@ -121,7 +121,7 @@ static void mkl_fa_online_softmax_chunk(
float * __restrict VKQ_accum,
int q0, int q_rows, int n_queries, int DV,
int chunk_size, int chunk_start,
int kvh_head, int gqa_ratio,
int kvh_head,
const sycl::half * mask_data, int64_t mask_head_stride,
int64_t mask_row_stride, int mask_n_heads,
float logit_softcap, int64_t wg_size) {
@ -473,7 +473,6 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor *
MKL_ACCUM(dequant_time_us, t_deq);
// --- Resolve mask pointers ---
const sycl::half * mask_data = nullptr;
int64_t mask_head_stride = 0;
int64_t mask_row_stride = 0;
int mask_n_heads = 0;
@ -547,7 +546,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor *
// 1. Pack all GQA Q heads into fp16 (full n_query_rows)
mkl_fa_pack_q_fp16(stream,
Q_head_f16_ptr, Q_batch,
n_queries, n_query_rows, DKQ,
n_queries, DKQ,
gqa_ratio, kvh_base_head,
q_scale, q_row_stride, q_head_stride, wg_size);
@ -605,7 +604,7 @@ void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor *
KQ_max_ptr, KQ_sum_ptr, VKQ_accum_ptr,
q0, q_rows, n_queries, DV,
this_chunk, chunk_start,
kvh_base_head, gqa_ratio,
kvh_base_head,
mask_batch, mask_head_stride,
mask_row_stride, mask_n_heads,
logit_softcap, wg_size);

20
ggml/src/ggml-sycl/ggml-sycl.cpp
View File

@ -921,16 +921,16 @@ ggml_backend_sycl_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft,
void * dev_ptr;
if (use_usm_system) {
GGML_SYCL_DEBUG("[SYCL] allocating %lu Bytes with USM system\n", size);
GGML_SYCL_DEBUG("[SYCL] allocating %zu Bytes with USM system\n", size);
dev_ptr = (void *)aligned_malloc_host(alignment, aligned_size);
if (!dev_ptr) {
GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on host\n", __func__, size);
GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on host\n", __func__, size);
return nullptr;
}
} else {
SYCL_CHECK(CHECK_TRY_ERROR(dev_ptr = (void *)ggml_sycl_malloc_device(size, *stream)));
if (!dev_ptr) {
GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on device\n", __func__, size);
GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on device\n", __func__, size);
return nullptr;
}
}
@ -1177,7 +1177,7 @@ ggml_backend_sycl_split_buffer_init_tensor(ggml_backend_buffer_t buffer,
SYCL_CHECK(CHECK_TRY_ERROR(buf = (char *)ggml_sycl_malloc_device(size, *stream)));
if (!buf) {
char err_buf[1024];
snprintf(err_buf, 1023, "%s: can't allocate %lu Bytes of memory on device\n", __func__, size);
snprintf(err_buf, 1023, "%s: can't allocate %zu Bytes of memory on device\n", __func__, size);
throw std::runtime_error(err_buf);
}
// set padding to 0 to avoid possible NaN values
@ -1651,7 +1651,7 @@ struct ggml_sycl_pool_leg : public ggml_sycl_pool {
SYCL_CHECK(CHECK_TRY_ERROR(ptr = (void *)ggml_sycl_malloc_device(look_ahead_size, *qptr)));
if (!ptr) {
GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on device/GPU\n", __func__, look_ahead_size);
GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on device/GPU\n", __func__, look_ahead_size);
return nullptr;
}
@ -1663,7 +1663,7 @@ struct ggml_sycl_pool_leg : public ggml_sycl_pool {
(uint32_t)(max_size/1024/1024), (uint32_t)(g_sycl_pool_size[id]/1024/1024), (uint32_t)(size/1024/1024));
#endif
// GGML_SYCL_DEBUG("ggml_sycl_pool_malloc_leg look_ahead_size=%lu, return %p\n", look_ahead_size, ptr);
// GGML_SYCL_DEBUG("ggml_sycl_pool_malloc_leg look_ahead_size=%zu, return %p\n", look_ahead_size, ptr);
return ptr;
}
@ -1843,7 +1843,7 @@ struct ggml_sycl_pool_host : public ggml_sycl_pool {
SYCL_CHECK(CHECK_TRY_ERROR(ptr = (void *) sycl::malloc_host(size, *qptr)));
if (!ptr) {
GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on host\n", __func__, size);
GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on host\n", __func__, size);
return nullptr;
}
pool_size += size;
@ -2779,9 +2779,9 @@ inline void ggml_sycl_op_mul_mat_sycl(
const float * src1_ddf1_i = src1->type == GGML_TYPE_F32 ? (const float *) src1_ddf_i : src1_ddq_as_f32.get();
{
#if GGML_SYCL_DNNL
const int64_t gemm_flops = (int64_t)row_diff * src1_ncols * ne10;
const bool use_mkl_direct = gemm_flops < 256 * 256 * 256;
#if GGML_SYCL_DNNL
if (g_ggml_sycl_enable_dnn && !use_mkl_direct) {
DnnlGemmWrapper::row_gemm(ctx, row_diff, src1_ncols, ne10, src0_ddf_i,
DnnlGemmWrapper::to_dt<float>(), src1_ddf1_i, DnnlGemmWrapper::to_dt<float>(),
@ -3518,7 +3518,9 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons
float * dst_ddf = static_cast<float *>(dst->data);
const sycl::half * src1_f16 = static_cast<const sycl::half *>(src1->data);
#if GGML_SYCL_DNNL
const size_t type_size_src0 = ggml_type_size(src0->type);
#endif
const size_t type_size_src1 = ggml_type_size(src1->type);
bool is_src0_cont_2 = ggml_is_contiguous_2(src0);
@ -3535,6 +3537,7 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons
scope_op_debug_print scope_dbg_print(__func__, "/to_fp16_nc_sycl", dst, /*num_src=*/2,
" : converting src1 to fp16");
#if GGML_SYCL_DNNL
// iterate tensor dims and find the slowest moving dim and stride
int last_dim=0;
int last_str=0;
@ -3554,7 +3557,6 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons
}
}
#if GGML_SYCL_DNNL
// oneDNN handles strided data and does not need overhead of ggml_get_to_fp16_nc_sycl
const int64_t ne_src1 = src1->nb[last_str] * src1->ne[last_dim] / type_size_src1;
src1_f16_alloc.alloc(ne_src1);

4
ggml/src/ggml-sycl/im2col.cpp
View File

@ -85,7 +85,7 @@ static void im2col_sycl(const float * x,
*/
stream->parallel_for(sycl::nd_range<3>(block_nums * sycl::range<3>(1, 1, MIN(IC_KH_KW, SYCL_IM2COL_BLOCK_SIZE)),
sycl::range<3>(1, 1, MIN(IC_KH_KW, SYCL_IM2COL_BLOCK_SIZE))),
[=](sycl::nd_item<3> item_ct1) {
[=](sycl::nd_item<3>) {
im2col_kernel(x, dst, IC, IW, IH, OH, OW, KW, KH, IC_IH_IW, IH_IW, N_OH, KH_KW, IC_KH_KW,
s0, s1, p0, p1, d0, d1);
});
@ -271,7 +271,7 @@ static void im2col_3d_sycl(const float * src,
*/
stream->parallel_for(sycl::nd_range<3>(block_nums * sycl::range<3>(1, 1, MIN(IC_KD_KH_KW, SYCL_IM2COL_BLOCK_SIZE)),
sycl::range<3>(1, 1, MIN(IC_KD_KH_KW, SYCL_IM2COL_BLOCK_SIZE))),
[=](sycl::nd_item<3> item_ct1) {
[=](sycl::nd_item<3>) {
im2col_3d_kernel(src, dst, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, OH_OW, KD_KH_KW,
ID_IH_IW, KH_KW, IH_IW, IC_ID_IH_IW, IC_KD_KH_KW, OW_KD_KH_KW,
OD_OH_OW_IC_KD_KH_KW, OH_OW_IC_KD_KH_KW, OW_IC_KD_KH_KW, N_OD_OH, OD_OH,

8
ggml/src/ggml-sycl/norm.cpp
View File

@ -7,9 +7,6 @@ static void norm_f32(const float* x, float* dst, const int ncols,
const int64_t dst_stride_col, const int64_t dst_stride_row, const int64_t dst_stride_channel, const int64_t dst_stride_sample,
const float eps, const sycl::nd_item<3>& item_ct1, sycl::float2* s_sum, int block_size) {
const int nrows = item_ct1.get_group_range(2);
const int nchannels = item_ct1.get_group_range(1);
const int nthreads = item_ct1.get_local_range(2);
const int sample = item_ct1.get_group(0);
const int channel = item_ct1.get_group(1);
@ -155,9 +152,6 @@ static void rms_norm_f32(const float* x, float* dst, const int ncols,
const float* mul = nullptr, const int64_t mul_stride_row = 0, const int64_t mul_stride_channel = 0,
const int64_t mul_stride_sample = 0, const int mul_nrows = 0, const int mul_nchannels = 0, const int mul_nsamples = 0) {
const int nrows = item_ct1.get_group_range(2);
const int nchannels = item_ct1.get_group_range(1);
const int sample = item_ct1.get_group(0);
const int channel = item_ct1.get_group(1);
const int row = item_ct1.get_group(2);
@ -225,8 +219,6 @@ static void l2_norm_f32(const float * x, float * dst, const int ncols,
const int64_t src_stride_sample, const int64_t dst_stride_col, const int64_t dst_stride_row,
const int64_t dst_stride_channel, const int64_t dst_stride_sample, const float eps,
const sycl::nd_item<3>& item_ct1, float* s_sum, const int block_size) {
const int nrows = item_ct1.get_group_range(2);
const int nchannels = item_ct1.get_group_range(1);
const int row = item_ct1.get_group(2);
const int channel = item_ct1.get_group(1);

2
ggml/src/ggml-sycl/set_rows.cpp
View File

@ -291,7 +291,7 @@ static void set_rows_sycl(
stream->parallel_for(
sycl::nd_range<1>(grid_size * block_size, block_size),
[=](sycl::nd_item<1> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] {
[=](sycl::nd_item<1> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
k_set_rows<TIn, TIdx, TOut>(
src0_d, src1_d, dst_d,
ne00, ne01, ne02,