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; 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 +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 }; + } else { + return sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::sub_group_size }; + } + } +}; + template static void lauch_kernel( dpct::dim3 group_range, @@ -881,7 +904,7 @@ static void lauch_kernel( static_cast>(group_range * local_range), static_cast>(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{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 #include "common.hpp" +#include "gla.hpp" template 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(); 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 { } 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.