diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index dc6cdd3df..661b526a9 100644 --- a/ggml/src/ggml-sycl/common.hpp +++ b/ggml/src/ggml-sycl/common.hpp @@ -214,6 +214,7 @@ inline dpct::err0 ggml_sycl_set_device(const int device) try { ////////////////////// struct optimize_feature { bool reorder=false; + bool onednn_optimized_gemm=false; }; struct sycl_device_info { @@ -637,6 +638,14 @@ constexpr size_t ceil_div(const size_t m, const size_t n) { bool gpu_has_xmx(sycl::device &dev); +#if GGML_SYCL_DNNL +// oneDNN builds JIT kernels only for some GPU architectures. On the rest it falls back to +// reference kernels, which are much slower than the SYCL kernels here. +inline bool ggml_sycl_dnnl_has_optimized_gemm(int device) { + return ggml_sycl_info().devices[device].opt_feature.onednn_optimized_gemm; +} +#endif + int ggml_sycl_get_env(const char *env_name, int default_val); template std::string debug_get_array_str(const std::string & prefix, const T array[N]) { diff --git a/ggml/src/ggml-sycl/fattn-onednn.cpp b/ggml/src/ggml-sycl/fattn-onednn.cpp index 4349363a3..6e030940a 100644 --- a/ggml/src/ggml-sycl/fattn-onednn.cpp +++ b/ggml/src/ggml-sycl/fattn-onednn.cpp @@ -34,6 +34,10 @@ bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst, bool use if (!g_ggml_sycl_fa_onednn) { return false; } + // without a fused SDPA kernel oneDNN falls back to the same matmul, so reuse the check + if (!ggml_sycl_dnnl_has_optimized_gemm(ggml_sycl_get_device())) { + return false; + } const ggml_tensor * Q = dst->src[0]; const ggml_tensor * K = dst->src[1]; const ggml_tensor * V = dst->src[2]; diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 6495ed433..e4e61e306 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -117,6 +117,41 @@ int g_ggml_sycl_enable_sparse_fa = 0; int g_ggml_sycl_debug_sparse_fa = 0; int g_ggml_sycl_sparse_fa_margin = 256; +#if GGML_SYCL_DNNL +// ask oneDNN which matmul it picks for a small f16 problem +static bool ggml_sycl_dnnl_detect_optimized_gemm(int device) { + using dt = dnnl::memory::data_type; + + try { + const dnnl::memory::dims dims = { 1, 64, 64 }; + const dnnl::memory::dims strides = { 64 * 64, 64, 1 }; + + dnnl::primitive_attr attr; + attr.set_scratchpad_mode(dnnl::scratchpad_mode::user); + + const auto & q = dpct::dev_mgr::instance().get_device(device).default_queue(); + const auto eng = dnnl::sycl_interop::make_engine(q.get_device(), q.get_context()); + const auto pd = dnnl::matmul::primitive_desc(eng, + dnnl::memory::desc(dims, dt::f16, strides), + dnnl::memory::desc(dims, dt::f16, strides), + dnnl::memory::desc(dims, dt::f32, strides), attr); + + const std::string impl = pd.impl_info_str(); + if (impl.find("ref") == std::string::npos) { + return true; + } + + GGML_LOG_WARN("%s: oneDNN has no optimized matmul for device %d (picks %s), using SYCL kernels\n", + __func__, device, impl.c_str()); + } catch (const std::exception & e) { + GGML_LOG_WARN("%s: oneDNN matmul probe failed on device %d (%s), using SYCL kernels\n", + __func__, device, e.what()); + } + + return false; +} +#endif + static ggml_sycl_device_info ggml_sycl_init() { GGML_SYCL_DEBUG("[SYCL] call ggml_sycl_init\n"); ggml_sycl_device_info info = {}; @@ -175,6 +210,9 @@ static ggml_sycl_device_info ggml_sycl_init() { 100 * prop.get_major_version() + 10 * prop.get_minor_version(); info.devices[i].nsm = prop.get_max_compute_units() / 16; //16: Number of Xe Cores info.devices[i].opt_feature.reorder = device.ext_oneapi_architecture_is(syclex::arch_category::intel_gpu); +#if GGML_SYCL_DNNL + info.devices[i].opt_feature.onednn_optimized_gemm = ggml_sycl_dnnl_detect_optimized_gemm(i); +#endif info.devices[i].smpbo = prop.get_local_mem_size(); info.devices[i].warp_size = WARP_SIZE; info.devices[i].usm_system_support = device.has(sycl::aspect::usm_system_allocations); @@ -2964,7 +3002,8 @@ inline void ggml_sycl_op_mul_mat_sycl( #if GGML_SYCL_DNNL && defined(GGML_SYCL_HAS_BF16) // Fast path for bf16 src0 - if (src0->type == GGML_TYPE_BF16 && g_ggml_sycl_enable_dnn && ggml_is_contiguous(src0) && + if (src0->type == GGML_TYPE_BF16 && g_ggml_sycl_enable_dnn && + ggml_sycl_dnnl_has_optimized_gemm(ggml_sycl_get_device()) && ggml_is_contiguous(src0) && row_diff == src0->ne[1]) { using bf16_t = sycl::ext::oneapi::bfloat16; ggml_sycl_pool_alloc src1_as_bf16(ctx.pool(), src1_ncols*ne10); @@ -3017,7 +3056,7 @@ inline void ggml_sycl_op_mul_mat_sycl( : src1_as_f16.get(); #if GGML_SYCL_DNNL - if (g_ggml_sycl_enable_dnn) { + if (g_ggml_sycl_enable_dnn && ggml_sycl_dnnl_has_optimized_gemm(ggml_sycl_get_device())) { DnnlGemmWrapper::row_gemm(ctx,row_diff, src1_ncols , ne10, src0_ptr, DnnlGemmWrapper::to_dt(), src1_ptr, DnnlGemmWrapper::to_dt(), dst_dd_i, DnnlGemmWrapper::to_dt(), stream); @@ -3061,7 +3100,7 @@ inline void ggml_sycl_op_mul_mat_sycl( #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 (g_ggml_sycl_enable_dnn && !use_mkl_direct) { + if (g_ggml_sycl_enable_dnn && !use_mkl_direct && ggml_sycl_dnnl_has_optimized_gemm(ggml_sycl_get_device())) { DnnlGemmWrapper::row_gemm(ctx, row_diff, src1_ncols, ne10, src0_ddf_i, DnnlGemmWrapper::to_dt(), src1_ddf1_i, DnnlGemmWrapper::to_dt(), dst_dd_i, DnnlGemmWrapper::to_dt(), stream); @@ -3892,7 +3931,7 @@ static void ggml_sycl_mul_mat_batched_sycl(ggml_backend_sycl_context & ctx, cons const int64_t r3 = ne13 / ne03; #if GGML_SYCL_DNNL - if (g_ggml_sycl_enable_dnn) { + if (g_ggml_sycl_enable_dnn && ggml_sycl_dnnl_has_optimized_gemm(ggml_sycl_get_device())) { int64_t str_a0 = nb00 / type_size_src0; int64_t str_a1 = nb01 / type_size_src0; int64_t str_a2 = nb02 / type_size_src0;