diff --git a/ggml-sycl.cpp b/ggml-sycl.cpp index d51f23b410595..edf8cc8445f9f 100644 --- a/ggml-sycl.cpp +++ b/ggml-sycl.cpp @@ -15272,8 +15272,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]; @@ -15282,16 +15282,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); @@ -15377,6 +15377,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__