From 10aef3680ff16dd7bab510cc5cf80782b5d73587 Mon Sep 17 00:00:00 2001 From: Aidan Date: Tue, 19 Mar 2024 16:07:41 +0000 Subject: [PATCH 1/3] Fix batched impl --- ggml-sycl.cpp | 20 +++++++++++--------- 1 file changed, 11 insertions(+), 9 deletions(-) diff --git a/ggml-sycl.cpp b/ggml-sycl.cpp index cc9ee0762b903..03949c9267d72 100644 --- a/ggml-sycl.cpp +++ b/ggml-sycl.cpp @@ -15274,8 +15274,8 @@ static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0, sycl_pool_alloc dst_f16; char * dst_t; - dpct::library_data_t cu_compute_type = dpct::library_data_t::real_float; - dpct::library_data_t cu_data_type = dpct::library_data_t::real_float; + dpct::library_data_t cu_compute_type = dpct::library_data_t::real_half; + dpct::library_data_t cu_data_type = dpct::library_data_t::real_half; // dst strides size_t nbd2 = dst->nb[2]; @@ -15284,16 +15284,16 @@ static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0, const sycl::half alpha_f16 = 1.0f; const sycl::half beta_f16 = 0.0f; - const float alpha_f32 = 1.0f; - const float beta_f32 = 0.0f; - - const void * alpha = &alpha_f32; - const void * beta = &beta_f32; + const void * alpha = &alpha_f16; + const void * beta = &beta_f16; // TODO: Renable (dst->op_params[0] =! GGML_PREC_DEFAULT) pathway - // oneMKL open source supports half, half, float, float: datatypes + // when oneMKL open source supports half, half, float, float: datatypes + + dst_t = (char *) dst_f16.alloc(ne_dst); - dst_t = (char *) dst_ddf; + nbd2 /= sizeof(float) / sizeof(sycl::half); + nbd3 /= sizeof(float) / sizeof(sycl::half); GGML_ASSERT(ne12 % ne02 == 0); GGML_ASSERT(ne13 % ne03 == 0); @@ -15379,6 +15379,8 @@ static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0, } #endif + const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16); + to_fp32_sycl(dst_f16.get(), dst_ddf, ne_dst, main_stream); } catch (sycl::exception const &exc) { std::cerr << exc.what() << "Exception caught at file:" << __FILE__ From d1cb8fedbd96687946eabe4a70e7a22c43a67898 Mon Sep 17 00:00:00 2001 From: Aidan Date: Thu, 21 Mar 2024 11:56:03 +0000 Subject: [PATCH 2/3] Maintain previous behaviour for igpu --- ggml-sycl.cpp | 37 ++++++++++++++++++++++++++++--------- 1 file changed, 28 insertions(+), 9 deletions(-) diff --git a/ggml-sycl.cpp b/ggml-sycl.cpp index 03949c9267d72..f4c0ad1d26df2 100644 --- a/ggml-sycl.cpp +++ b/ggml-sycl.cpp @@ -15246,6 +15246,9 @@ static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0, SYCL_CHECK(ggml_sycl_set_device(g_main_device)); dpct::queue_ptr main_stream = g_syclStreams[g_main_device][0]; + bool no_mixed_dtypes = main_stream->get_backend() == sycl::backend::ext_oneapi_cuda || + main_stream->get_backend() == sycl::backend::ext_oneapi_hip; + SYCL_CHECK( CHECK_TRY_ERROR(g_sycl_handles[g_main_device] = main_stream)); @@ -15274,26 +15277,40 @@ static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0, sycl_pool_alloc dst_f16; char * dst_t; - dpct::library_data_t cu_compute_type = dpct::library_data_t::real_half; - dpct::library_data_t cu_data_type = dpct::library_data_t::real_half; + dpct::library_data_t cu_compute_type = dpct::library_data_t::real_float; + dpct::library_data_t cu_data_type = dpct::library_data_t::real_float; + if (no_mixed_dtypes) { + cu_compute_type = dpct::library_data_t::real_half; + cu_data_type = dpct::library_data_t::real_half; + } // dst strides size_t nbd2 = dst->nb[2]; size_t nbd3 = dst->nb[3]; + const float alpha_f32 = 1.0f; + const float beta_f32 = 0.0f; + const sycl::half alpha_f16 = 1.0f; const sycl::half beta_f16 = 0.0f; - const void * alpha = &alpha_f16; - const void * beta = &beta_f16; + const void * alpha = &alpha_f32; + const void * beta = &beta_f32; + if (no_mixed_dtypes) { + alpha = &alpha_f16; + beta = &beta_f16; + } // TODO: Renable (dst->op_params[0] =! GGML_PREC_DEFAULT) pathway // when oneMKL open source supports half, half, float, float: datatypes - dst_t = (char *) dst_f16.alloc(ne_dst); + dst_t = (char *) dst_ddf; + if (no_mixed_dtypes) { + dst_t = (char *) dst_f16.alloc(ne_dst); - nbd2 /= sizeof(float) / sizeof(sycl::half); - nbd3 /= sizeof(float) / sizeof(sycl::half); + nbd2 /= sizeof(float) / sizeof(sycl::half); + nbd3 /= sizeof(float) / sizeof(sycl::half); + } GGML_ASSERT(ne12 % ne02 == 0); GGML_ASSERT(ne13 % ne03 == 0); @@ -15379,8 +15396,10 @@ static void ggml_sycl_mul_mat_batched_sycl(const ggml_tensor *src0, } #endif - const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16); - to_fp32_sycl(dst_f16.get(), dst_ddf, ne_dst, main_stream); + if (no_mixed_dtypes) { + const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16); + to_fp32_sycl(dst_f16.get(), dst_ddf, ne_dst, main_stream); + } } catch (sycl::exception const &exc) { std::cerr << exc.what() << "Exception caught at file:" << __FILE__ From cec6481ae029e32ea78bb67b0f961930595c8a87 Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Wed, 27 Mar 2024 11:23:04 +0530 Subject: [PATCH 3/3] retrigger CI