diff --git a/ggml/src/ggml-sycl/fattn-vec.hpp b/ggml/src/ggml-sycl/fattn-vec.hpp index 9ec88c287..5e40c88fb 100644 --- a/ggml/src/ggml-sycl/fattn-vec.hpp +++ b/ggml/src/ggml-sycl/fattn-vec.hpp @@ -589,10 +589,10 @@ void ggml_sycl_flash_attn_ext_vec_case_impl(ggml_backend_sycl_context & ctx, ggm const bool need_f16_V = type_V == GGML_TYPE_F16; constexpr size_t nbytes_shared = 0; + const auto arch = ggml_sycl_info().devices[ggml_sycl_get_device()].hw_info.arch; // D=512 does not fit the default register file; it spills up to 343 bytes per thread, against at most 57 for D <= 256. This kernel is decode only, so thread occupancy is not the limit. It is 1.9x faster at every KV depth on Battlemage. constexpr bool use_large_grf = D >= 512; - const auto arch = ggml_sycl_info().devices[ctx.device].hw_info.arch; const int nthreads = ggml_sycl_fattn_vec_get_nthreads_device(arch); if constexpr (D <= 256) { if (nthreads == 256) { diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index c6aff3955..b4d200526 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -792,6 +792,14 @@ static bool ggml_sycl_is_l0_discrete_gpu(int device) { } #endif +static void memcpy_host_forward(sycl::queue &q_dst, sycl::queue &q_src, void *ptr_dst, + const void *ptr_src, size_t size) { + char *host_buf = (char *)malloc(size); + q_src.memcpy(host_buf, (const char *)ptr_src, size).wait(); + q_dst.memcpy((char *)ptr_dst, host_buf, size).wait(); + free(host_buf); +} + static void dev2dev_memcpy(int device_dst, sycl::queue &q_dst, int device_src, sycl::queue &q_src, void *ptr_dst, const void *ptr_src, size_t size) { @@ -835,10 +843,7 @@ static void dev2dev_memcpy(int device_dst, sycl::queue &q_dst, int device_src, s } else { GGML_SYCL_DEBUG("[SYCL] dev2dev memcpy by host forward for SYCL/L0 fallback\n"); } - char *host_buf = (char *)malloc(size); - q_src.memcpy(host_buf, (const char *)ptr_src, size).wait(); - q_dst.memcpy((char *)ptr_dst, host_buf, size).wait(); - free(host_buf); + memcpy_host_forward(q_dst, q_src, ptr_dst, ptr_src, size); } static bool