mirror of
https://github.com/ggml-org/whisper.cpp.git
synced 2026-10-21 09:41:29 +02:00
fix the issue in mixed different model GPUs in FA (llama/29071)
* fix mixed different model GPUs issue * Update ggml/src/ggml-sycl/ggml-sycl.cpp Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
This commit is contained in:
co-authored by
Georgi Gerganov
parent
24be1aca17
commit
b35b457a14
@@ -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;
|
const bool need_f16_V = type_V == GGML_TYPE_F16;
|
||||||
constexpr size_t nbytes_shared = 0;
|
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.
|
// 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;
|
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);
|
const int nthreads = ggml_sycl_fattn_vec_get_nthreads_device(arch);
|
||||||
if constexpr (D <= 256) {
|
if constexpr (D <= 256) {
|
||||||
if (nthreads == 256) {
|
if (nthreads == 256) {
|
||||||
|
|||||||
@@ -792,6 +792,14 @@ static bool ggml_sycl_is_l0_discrete_gpu(int device) {
|
|||||||
}
|
}
|
||||||
#endif
|
#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,
|
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) {
|
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 {
|
} else {
|
||||||
GGML_SYCL_DEBUG("[SYCL] dev2dev memcpy by host forward for SYCL/L0 fallback\n");
|
GGML_SYCL_DEBUG("[SYCL] dev2dev memcpy by host forward for SYCL/L0 fallback\n");
|
||||||
}
|
}
|
||||||
char *host_buf = (char *)malloc(size);
|
memcpy_host_forward(q_dst, q_src, ptr_dst, ptr_src, 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 bool
|
static bool
|
||||||
|
|||||||
Reference in New Issue
Block a user