Repository: ggml-org/llama.cpp
Pull request: sycl: Remove not needed copy f16->f32 for dnnl mul mat

PR proposes when `GGML_SYCL_F16=ON` to allow oneDNN to handle conversion and outputting of mul_mat into f32 and enabling fpmathmode to `f16`.

The current approach uses the memory pool to pass a f16 `dst` for the oneDNN matmul and a `cpy` from f16 to the actual f32 output `dst_dd_i`
Example of performance difference observed: 

### Lunar Lake

current approach:
model | size | params | backend | ngl | threads | test | t/s
-- | -- | -- | -- | -- | -- | -- | --
qwen2 1.5B Q4_0 | 1013.62 MiB | 1.78 B | SYCL | 99 | 8 | pp512 | 1475.42 ± 43.91
qwen2 1.5B Q4_0 | 1013.62 MiB | 1.78 B | SYCL | 99 | 8 | tg128 | 37.46 ± 0.38

proposed changes:
model | size | params | backend | ngl | threads | test | t/s
-- | -- | -- | -- | -- | -- | -- | --
qwen2 1.5B Q4_0 | 1013.62 MiB | 1.78 B | SYCL | 99 | 8 | pp512 | 1566.90 ± 26.33
qwen2 1.5B Q4_0 | 1013.62 MiB | 1.78 B | SYCL | 99 | 8 | tg128 | 37.66 ± 0.49

### Battlemage

current approach:
| model | size | params | backend | ngl | test | t/s |
| --- | --- | --- | --- | --- | --- | --- |
| qwen2 1.5B Q4_0 | 1013.62 MiB | 1.78 B | SYCL | 99 | pp512 | 7424.59 ± 15.58 |
| qwen2 1.5B Q4_0 | 1013.62 MiB | 1.78 B | SYCL | 99 | tg128 | 99.61 ± 2.09 |

proposed changes:
| model | size | params | backend | ngl | test | t/s |
| --- | --- | --- | --- | --- | --- | --- |
| qwen2 1.5B Q4_0 | 1013.62 MiB | 1.78 B | SYCL | 99 | pp512 | 8148.99 ± 35.10 |
| qwen2 1.5B Q4_0 | 1013.62 MiB | 1.78 B | SYCL | 99 | tg128 | 100.62 ± 2.05 |


Changed files:
filename: ggml/src/ggml-sycl/gemm.hpp
patch: @@ -65,6 +65,9 @@ class DnnlGemmWrapper {
 
         dnnl::primitive_attr primitive_attr;
         primitive_attr.set_scratchpad_mode(dnnl::scratchpad_mode::user);
+#ifdef GGML_SYCL_F16
+        primitive_attr.set_fpmath_mode(dnnl::fpmath_mode::f16);
+#endif
 
         auto a_mem = dnnl::memory(a_in_md, eng, const_cast<void*>(a));
         auto b_mem = dnnl::memory(b_in_md, eng, const_cast<void*>(b));

filename: ggml/src/ggml-sycl/ggml-sycl.cpp
patch: @@ -2127,21 +2127,18 @@ inline void ggml_sycl_op_mul_mat_sycl(
         const sycl::half *src1_ptr = src1->type == GGML_TYPE_F16
                 ? (const sycl::half *)src1->data + src1_padded_row_size
                                          : src1_as_f16.get();
-        ggml_sycl_pool_alloc<sycl::half> dst_f16(ctx.pool(), row_diff * src1_ncols);
 
 #if GGML_SYCL_DNNL
         if (!g_ggml_sycl_disable_dnn) {
             DnnlGemmWrapper::row_gemm(ctx, src1_ncols, row_diff, ne10, src1_ptr,
                                       DnnlGemmWrapper::to_dt<sycl::half>(), src0_ptr, DnnlGemmWrapper::to_dt<sycl::half>(),
-                                      dst_f16.get(), DnnlGemmWrapper::to_dt<sycl::half>(), stream);
-            scope_op_debug_print scope_dbg_print(__func__, "/to_fp32_sycl", dst, /*num_src=*/2,
-                                                 " : converting dst to fp32");
-            const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16, dst);
-            to_fp32_sycl(dst_f16.get(), dst_dd_i, row_diff* src1_ncols, stream);
+                                      dst_dd_i, DnnlGemmWrapper::to_dt<float>(), stream);
         }
         else
 #endif
         {
+            ggml_sycl_pool_alloc<sycl::half> dst_f16(ctx.pool(), row_diff * src1_ncols);
+
             const sycl::half alpha_f16 = 1.0f;
             const sycl::half beta_f16  = 0.0f;
             SYCL_CHECK(CHECK_TRY_ERROR(dpct::gemm(