Commit dd002b91f for llama.cpp
commit dd002b91f583af0b80140f9a140dc04856bfc2e0
Author: Michael Taylor <162068037+mctylr-gh@users.noreply.github.com>
Date: Tue Oct 20 01:47:25 2026 -0300
sycl: fix compiler warnings - missing prototypes, dead code, extra semicolons (#30271)
diff --git a/cmake/common.cmake b/cmake/common.cmake
index bcf403e0e..ee8d5bf7b 100644
--- a/cmake/common.cmake
+++ b/cmake/common.cmake
@@ -2,7 +2,7 @@ include("ggml/cmake/common.cmake")
function(llama_add_compile_flags)
if (LLAMA_FATAL_WARNINGS)
- if (CMAKE_CXX_COMPILER_ID MATCHES "GNU" OR CMAKE_CXX_COMPILER_ID MATCHES "Clang")
+ if (CMAKE_CXX_COMPILER_ID MATCHES "GNU" OR CMAKE_CXX_COMPILER_ID MATCHES "Clang" OR CMAKE_CXX_COMPILER_ID MATCHES "IntelLLVM")
list(APPEND C_FLAGS -Werror)
list(APPEND CXX_FLAGS -Werror)
elseif (CMAKE_CXX_COMPILER_ID STREQUAL "MSVC")
diff --git a/ggml/cmake/common.cmake b/ggml/cmake/common.cmake
index f3610298f..42b65a580 100644
--- a/ggml/cmake/common.cmake
+++ b/ggml/cmake/common.cmake
@@ -2,10 +2,16 @@ function(ggml_get_flags CCID CCVER)
set(C_FLAGS "")
set(CXX_FLAGS "")
- if (CCID MATCHES "Clang")
+ if (CCID MATCHES "Clang" OR CCID MATCHES "IntelLLVM")
set(C_FLAGS -Wunreachable-code-break -Wunreachable-code-return)
set(CXX_FLAGS -Wunreachable-code-break -Wunreachable-code-return -Wmissing-prototypes -Wextra-semi)
+ if (CCID STREQUAL "IntelLLVM")
+ # for Intel DPC++ and LLVM SYCL, using debug does not disable optimization
+ list(APPEND C_FLAGS -Rno-debug-disables-optimization)
+ list(APPEND CXX_FLAGS -Rno-debug-disables-optimization)
+ endif()
+
if (
(CCID STREQUAL "Clang" AND CCVER VERSION_GREATER_EQUAL 3.8.0) OR
(CCID STREQUAL "AppleClang" AND CCVER VERSION_GREATER_EQUAL 7.3.0)
diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt
index dfdb4746d..71c76737f 100644
--- a/ggml/src/ggml-sycl/CMakeLists.txt
+++ b/ggml/src/ggml-sycl/CMakeLists.txt
@@ -44,7 +44,7 @@ if (WIN32)
if(DEFINED ENV{LEVEL_ZERO_V1_SDK_PATH})
set(LEVEL_ZERO_V1_SDK_PATH $ENV{LEVEL_ZERO_V1_SDK_PATH})
if(EXISTS "${LEVEL_ZERO_V1_SDK_PATH}")
- target_include_directories(ggml-sycl PRIVATE "${LEVEL_ZERO_V1_SDK_PATH}/include")
+ target_include_directories(ggml-sycl SYSTEM PRIVATE "${LEVEL_ZERO_V1_SDK_PATH}/include")
set(LEVEL_ZERO_V1_SDK_LIB_PATH "${LEVEL_ZERO_V1_SDK_PATH}/lib")
else()
message(WARNING "LEVEL_ZERO_V1_SDK_PATH set but folder not found: ${LEVEL_ZERO_V1_SDK_PATH}")
diff --git a/ggml/src/ggml-sycl/concat.cpp b/ggml/src/ggml-sycl/concat.cpp
index bd5f3b2ce..d4c257cf7 100644
--- a/ggml/src/ggml-sycl/concat.cpp
+++ b/ggml/src/ggml-sycl/concat.cpp
@@ -504,6 +504,5 @@ void ggml_sycl_op_concat(ggml_backend_sycl_context & ctx, ggml_tensor *dst) {
default:
fprintf(stderr, "%s: unsupported types: dst: %s\n", __func__, ggml_type_name(dst->type));
GGML_ASSERT(false);
- break;
}
}
diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp
index 45c60e8e5..092967b6f 100644
--- a/ggml/src/ggml-sycl/convert.cpp
+++ b/ggml/src/ggml-sycl/convert.cpp
@@ -757,7 +757,6 @@ to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) {
#endif
default:
GGML_ABORT("fatal error: unsupport data type=%s\n", ggml_type_name(type));
- return nullptr;
}
}
@@ -861,7 +860,6 @@ to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) {
#endif
default:
GGML_ABORT("fatal error: unsupport data type=%s\n", ggml_type_name(type));
- return nullptr;
}
}
@@ -877,7 +875,6 @@ to_bf16_sycl_t ggml_get_to_bf16_sycl(ggml_type type, ggml_tensor * /*dst*/) {
return convert_unary_sycl<sycl::ext::oneapi::bfloat16>;
default:
GGML_ABORT("fatal error: unsupport data type=%s\n", ggml_type_name(type));
- return nullptr;
}
}
#endif
diff --git a/ggml/src/ggml-sycl/count-equal.cpp b/ggml/src/ggml-sycl/count-equal.cpp
index 4580354cd..20126652b 100644
--- a/ggml/src/ggml-sycl/count-equal.cpp
+++ b/ggml/src/ggml-sycl/count-equal.cpp
@@ -74,6 +74,5 @@ void ggml_sycl_count_equal(ggml_backend_sycl_context &ctx, ggml_tensor *dst) {
} break;
default:
GGML_ASSERT(false);
- break;
}
}
diff --git a/ggml/src/ggml-sycl/dpct/helper.hpp b/ggml/src/ggml-sycl/dpct/helper.hpp
index 85af4cab6..fd3b9655a 100644
--- a/ggml/src/ggml-sycl/dpct/helper.hpp
+++ b/ggml/src/ggml-sycl/dpct/helper.hpp
@@ -3159,7 +3159,7 @@ namespace dpct
_local_mem_size = local_mem_size;
- };
+ }
static inline std::mutex kernel_function_ptr_map_mutex;
public:
diff --git a/ggml/src/ggml-sycl/element_wise.cpp b/ggml/src/ggml-sycl/element_wise.cpp
index 0d263aa76..c7fe70ac3 100644
--- a/ggml/src/ggml-sycl/element_wise.cpp
+++ b/ggml/src/ggml-sycl/element_wise.cpp
@@ -1169,7 +1169,7 @@ static void swiglu_oai_sycl(const T * x,
});
}
-void ggml_sycl_op_swiglu_oai(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+static void ggml_sycl_op_swiglu_oai(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
const ggml_tensor * src1 = dst->src[1];
void * src0_d = src0->data;
diff --git a/ggml/src/ggml-sycl/fattn-common.hpp b/ggml/src/ggml-sycl/fattn-common.hpp
index 2e8ac6af3..c37326336 100644
--- a/ggml/src/ggml-sycl/fattn-common.hpp
+++ b/ggml/src/ggml-sycl/fattn-common.hpp
@@ -832,6 +832,29 @@ static void flash_attn_combine_results(const float * __restrict__ VKQ_parts,
dst[tid] = VKQ_numerator / VKQ_denominator;
}
+// Carries the sub-group size (and grf_size_automatic when use_large_grf) as
+// kernel properties via get(properties_tag), in place of the
+// reqd_sub_group_size attribute and the deprecated parallel_for overload.
+template <typename KernelFunc, uint32_t Size, bool UseGrf>
+struct props_kernel {
+ KernelFunc kernel;
+
+ void operator()(sycl::nd_item<3> item) const {
+ kernel(item);
+ }
+
+ auto get(sycl::ext::oneapi::experimental::properties_tag) const {
+ if constexpr (UseGrf) {
+ return sycl::ext::oneapi::experimental::properties{
+ sycl::ext::intel::experimental::grf_size_automatic,
+ sycl::ext::oneapi::experimental::sub_group_size<Size> };
+ } else {
+ return sycl::ext::oneapi::experimental::properties{
+ sycl::ext::oneapi::experimental::sub_group_size<Size> };
+ }
+ }
+};
+
template <fattn_kernel_t fattn_kernel, int warp_size, bool use_large_grf = false>
static void lauch_kernel(
dpct::dim3 group_range,
@@ -881,7 +904,7 @@ static void lauch_kernel(
static_cast<sycl::range<3>>(group_range * local_range),
static_cast<sycl::range<3>>(local_range));
- const auto kernel = [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(warp_size)]] {
+ const auto kernel = [=](sycl::nd_item<3> item_ct1) {
GGML_UNUSED(item_ct1);
fattn_kernel(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
max_bias, m0, m1, n_head_log2, logit_softcap, ne00,
@@ -892,14 +915,7 @@ static void lauch_kernel(
q->submit([&](sycl::handler &cgh) {
// grf_size_automatic lets the compiler use the large register file when the kernel needs it. grf_size<256> would say the same but is undefined on devices that do not have it.
- if constexpr (use_large_grf) {
- cgh.parallel_for(
- rng,
- sycl::ext::oneapi::experimental::properties{ sycl::ext::intel::experimental::grf_size_automatic },
- kernel);
- } else {
- cgh.parallel_for(rng, kernel);
- }
+ cgh.parallel_for(rng, props_kernel<decltype(kernel), warp_size, use_large_grf>{kernel});
});
}
diff --git a/ggml/src/ggml-sycl/fattn-tile.cpp b/ggml/src/ggml-sycl/fattn-tile.cpp
index 9449d7578..db030268e 100644
--- a/ggml/src/ggml-sycl/fattn-tile.cpp
+++ b/ggml/src/ggml-sycl/fattn-tile.cpp
@@ -54,6 +54,6 @@ void ggml_sycl_flash_attn_ext_tile(ggml_backend_sycl_context & ctx, ggml_tensor
} break;
default: {
GGML_ABORT("Unsupported head size");
- } break;
+ }
}
}
diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp
index 9c17daf86..a3a320900 100644
--- a/ggml/src/ggml-sycl/getrows.cpp
+++ b/ggml/src/ggml-sycl/getrows.cpp
@@ -652,6 +652,5 @@ void ggml_sycl_op_get_rows_back(ggml_backend_sycl_context & ctx, ggml_tensor * d
break;
default:
GGML_ABORT("%s: unsupported src0 type: %s\n", __func__, ggml_type_name(src0->type));
- break;
}
}
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 3cc565d6b..cbb2188c5 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -388,14 +388,14 @@ static ze_result_t init_zes() {
return res;
}
-ze_result_t get_zes_init_res() {
+static ze_result_t get_zes_init_res() {
static ze_result_t zes_init_res = init_zes();
GGML_SYCL_DEBUG("[SYCL] call %s: zesInit result: %d\n", __func__, (int) zes_init_res);
return zes_init_res;
}
#endif
-void initialize_sycl_begining() {
+static void initialize_sycl_begining() {
#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
//must be called in initialization stage, before any other Level Zero API calls
GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__);
@@ -1173,7 +1173,7 @@ static size_t ggml_backend_sycl_buffer_type_get_alignment(ggml_backend_buffer_ty
GGML_UNUSED(buft);
}
-bool is_bmg_g31_arch(int device) {
+static bool is_bmg_g31_arch(int device) {
return ggml_sycl_info().devices[device].hw_info.arch == gpu_arch::intel_gpu_bmg_g31;
}
@@ -6326,7 +6326,7 @@ catch (sycl::exception const &exc) {
std::exit(1);
}
-bool sycl_get_mem_info(int device, size_t * free, size_t * total) {
+static bool sycl_get_mem_info(int device, size_t * free, size_t * total) {
GGML_SYCL_DEBUG("[SYCL] [%s] g_ggml_sycl_get_mem_api=%d\n",
__func__, g_ggml_sycl_get_mem_api);
@@ -6952,7 +6952,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
default:
return false;
}
- break;
case GGML_OP_MUL_MAT:
case GGML_OP_MUL_MAT_ID:
{
@@ -7045,7 +7044,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
(op->src[1]->type == GGML_TYPE_I64 || op->src[1]->type == GGML_TYPE_I32);
return res;
}
- break;
case GGML_OP_DSV4_HC_PRE:
return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 &&
op->type == GGML_TYPE_F32;
diff --git a/ggml/src/ggml-sycl/gla.cpp b/ggml/src/ggml-sycl/gla.cpp
index 879184fdd..379483427 100644
--- a/ggml/src/ggml-sycl/gla.cpp
+++ b/ggml/src/ggml-sycl/gla.cpp
@@ -1,6 +1,7 @@
#include <sycl/sycl.hpp>
#include "common.hpp"
+#include "gla.hpp"
template <u_int HEAD_SIZE>
static void gated_linear_attn_f32_kernel(const dpct::queue_ptr stream, u_int B, u_int T, u_int C, u_int H, float scale,
diff --git a/ggml/src/ggml-sycl/mem.cpp b/ggml/src/ggml-sycl/mem.cpp
index ad5bfe0ff..8a1c3262c 100644
--- a/ggml/src/ggml-sycl/mem.cpp
+++ b/ggml/src/ggml-sycl/mem.cpp
@@ -27,7 +27,7 @@ const char * mem_api_int2str(int mem_api) {
/*
* Depend on to call zesInit(0) before any other Level Zero API calls, otherwise the Level Zero API calls may fail.
*/
-bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
+static bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using Level Zero API.\n", __func__);
free_bytes = 0;
@@ -94,7 +94,7 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot
}
#endif
-bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
+static bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using SYCL API.\n", __func__);
total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
diff --git a/ggml/src/ggml-sycl/topk-moe.cpp b/ggml/src/ggml-sycl/topk-moe.cpp
index 78574c4b5..3dd92d261 100644
--- a/ggml/src/ggml-sycl/topk-moe.cpp
+++ b/ggml/src/ggml-sycl/topk-moe.cpp
@@ -296,7 +296,6 @@ static void ggml_sycl_op_topk_moe(ggml_backend_sycl_context & ctx,
break;
default:
GGML_ASSERT(false && "fatal error");
- break;
}
}
diff --git a/ggml/src/ggml-sycl/upscale.cpp b/ggml/src/ggml-sycl/upscale.cpp
index e42cb419d..5f1c16be2 100644
--- a/ggml/src/ggml-sycl/upscale.cpp
+++ b/ggml/src/ggml-sycl/upscale.cpp
@@ -183,8 +183,8 @@ static void upscale_f32_bilinear_antialias(const float * src0,
}
namespace bicubic_interpolation {
-static float weight1(float x, const float &a) { return ((a + 2) * x - (a + 3)) * x * x + 1; };
-static float weight2(float x, const float &a) { return ((a * x - 5 * a) * x + 8 * a) * x - 4 * a; };
+static float weight1(float x, const float &a) { return ((a + 2) * x - (a + 3)) * x * x + 1; }
+static float weight2(float x, const float &a) { return ((a * x - 5 * a) * x + 8 * a) * x - 4 * a; }
static float bicubic(float p0, float p1, float p2, float p3, float x, float a) {
const float w0 = weight2(x + 1, a);
@@ -192,7 +192,7 @@ static float bicubic(float p0, float p1, float p2, float p3, float x, float a) {
const float w2 = weight1(1 - x, a);
const float w3 = weight2(2 - x, a);
return p0 * w0 + p1 * w1 + p2 * w2 + p3 * w3;
-};
+}
}
@@ -358,7 +358,7 @@ static void upscale_f32_bicubic_sycl(const float * x,
}
}
-void ggml_sycl_op_upscale(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
+static void ggml_sycl_op_upscale(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
const ggml_tensor * src0 = dst->src[0];
const float * src0_d = (const float *)src0->data;
float * dst_d = (float *)dst->data;
diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp
index 57360a004..f547b20e6 100644
--- a/ggml/src/ggml-sycl/vecdotq.hpp
+++ b/ggml/src/ggml-sycl/vecdotq.hpp
@@ -442,7 +442,7 @@ template <> struct reorder_vec_dot_q_sycl<GGML_TYPE_Q4_0> {
}
return vec_dot_q4_0_q8_1_impl(v, u, d, *q8_1_ds);
- };
+ }
};
// Load four contiguous dwords per operand instead of loading each value separately.