From 75ab2d06f52477ba1957b764873a3f18c4359639 Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Sun, 28 Jan 2024 21:26:23 +0530 Subject: [PATCH] ggml : add unified SYCL backend for Intel GPUs (llama/2690) * first update for migration * update init_cublas * add debug functio, commit all help code * step 1 * step 2 * step3 add fp16, slower 31->28 * add GGML_LIST_DEVICE function * step 5 format device and print * step6, enhance error check, remove CUDA macro, enhance device id to fix none-zero id issue * support main device is non-zero * step7 add debug for code path, rm log * step 8, rename all macro & func from cuda by sycl * fix error of select non-zero device, format device list * ren ggml-sycl.hpp -> ggml-sycl.h * clear CMAKE to rm unused lib and options * correct queue: rm dtct:get_queue * add print tensor function to debug * fix error: wrong result in 658746bb26702e50f2c59c0e4ada8e9da6010481 * summary dpct definition in one header file to replace folder:dpct * refactor device log * mv dpct definition from folder dpct to ggml-sycl.h * update readme, refactor build script * fix build with sycl * set nthread=1 when sycl, increase performance * add run script, comment debug code * add ls-sycl-device tool * add ls-sycl-device, rm unused files * rm rear space * dos2unix * Update README_sycl.md * fix return type * remove sycl version from include path * restore rm code to fix hang issue * add syc and link for sycl readme * rm original sycl code before refactor * fix code err * add know issue for pvc hang issue * enable SYCL_F16 support * align pr4766 * check for sycl blas, better performance * cleanup 1 * remove extra endif * add build&run script, clean CMakefile, update guide by review comments * rename macro to intel hardware * editor config format * format fixes * format fixes * editor format fix * Remove unused headers * skip build sycl tool for other code path * replace tab by space * fix blas matmul function * fix mac build * restore hip dependency * fix conflict * ren as review comments * mv internal function to .cpp file * export funciton print_sycl_devices(), mv class dpct definition to source file * update CI/action for sycl code, fix CI error of repeat/dup * fix action ID format issue * rm unused strategy * enable llama_f16 in ci * fix conflict * fix build break on MacOS, due to CI of MacOS depend on external ggml, instead of internal ggml * fix ci cases for unsupported data type * revert unrelated changed in cuda cmake remove useless nommq fix typo of GGML_USE_CLBLAS_SYCL * revert hip cmake changes * fix indent * add prefix in func name * revert no mmq * rm cpu blas duplicate * fix no_new_line * fix src1->type==F16 bug. * pass batch offset for F16 src1 * fix batch error * fix wrong code * revert sycl checking in test-sampling * pass void as arguments of ggml_backend_sycl_print_sycl_devices * remove extra blank line in test-sampling * revert setting n_threads in sycl * implement std::isinf for icpx with fast math. * Update ci/run.sh Co-authored-by: Georgi Gerganov * Update examples/sycl/run-llama2.sh Co-authored-by: Georgi Gerganov * Update examples/sycl/run-llama2.sh Co-authored-by: Georgi Gerganov * Update CMakeLists.txt Co-authored-by: Georgi Gerganov * Update CMakeLists.txt Co-authored-by: Georgi Gerganov * Update CMakeLists.txt Co-authored-by: Georgi Gerganov * Update CMakeLists.txt Co-authored-by: Georgi Gerganov * add copyright and MIT license declare * update the cmd example --------- Co-authored-by: jianyuzh Co-authored-by: luoyu-intel Co-authored-by: Meng, Hengyu Co-authored-by: Georgi Gerganov --- ggml-backend.c | 5 +++++ ggml.c | 22 ++++++++++++++++++++-- ggml.h | 1 + 3 files changed, 26 insertions(+), 2 deletions(-) diff --git a/ggml-backend.c b/ggml-backend.c index 3fff5fc..897a4cb 100644 --- a/ggml-backend.c +++ b/ggml-backend.c @@ -339,6 +339,11 @@ GGML_CALL static void ggml_backend_registry_init(void) { ggml_backend_cuda_reg_devices(); #endif +#ifdef GGML_USE_SYCL + extern void ggml_backend_sycl_reg_devices(void); + ggml_backend_sycl_reg_devices(); +#endif + #ifdef GGML_USE_METAL extern GGML_CALL ggml_backend_t ggml_backend_reg_metal_init(const char * params, void * user_data); extern GGML_CALL ggml_backend_buffer_type_t ggml_backend_metal_buffer_type(void); diff --git a/ggml.c b/ggml.c index cf433e9..8236ff5 100644 --- a/ggml.c +++ b/ggml.c @@ -248,6 +248,8 @@ inline static void * ggml_aligned_malloc(size_t size) { #include "ggml-cuda.h" #elif defined(GGML_USE_CLBLAST) #include "ggml-opencl.h" +#elif defined(GGML_USE_SYCL) +#include "ggml-sycl.h" #endif // floating point type used to accumulate sums @@ -2293,6 +2295,8 @@ struct ggml_context * ggml_init(struct ggml_init_params params) { ggml_init_cublas(); #elif defined(GGML_USE_CLBLAST) ggml_cl_init(); +#elif defined(GGML_USE_SYCL) + ggml_init_sycl(); #endif ggml_setup_op_has_task_pass(); @@ -14701,6 +14705,12 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm GGML_ASSERT(tensor->src[1] == NULL || tensor->src[1]->backend == GGML_BACKEND_CPU); #endif // GGML_USE_CUBLAS +#ifdef GGML_USE_SYCL + bool skip_cpu = ggml_sycl_compute_forward(params, tensor); + if (skip_cpu) { + return; + } +#endif // GGML_USE_SYCL switch (tensor->op) { case GGML_OP_DUP: { @@ -20280,7 +20290,7 @@ int ggml_cpu_has_wasm_simd(void) { } int ggml_cpu_has_blas(void) { -#if defined(GGML_USE_ACCELERATE) || defined(GGML_USE_OPENBLAS) || defined(GGML_USE_CUBLAS) || defined(GGML_USE_CLBLAST) +#if defined(GGML_USE_ACCELERATE) || defined(GGML_USE_OPENBLAS) || defined(GGML_USE_CUBLAS) || defined(GGML_USE_CLBLAST) || defined(GGML_USE_SYCL) return 1; #else return 0; @@ -20303,8 +20313,16 @@ int ggml_cpu_has_clblast(void) { #endif } +int ggml_cpu_has_sycl(void) { +#if defined(GGML_USE_SYCL) + return 1; +#else + return 0; +#endif +} + int ggml_cpu_has_gpublas(void) { - return ggml_cpu_has_cublas() || ggml_cpu_has_clblast(); + return ggml_cpu_has_cublas() || ggml_cpu_has_clblast() || ggml_cpu_has_sycl(); } int ggml_cpu_has_sse3(void) { diff --git a/ggml.h b/ggml.h index 1c49762..3d8d6f2 100644 --- a/ggml.h +++ b/ggml.h @@ -2266,6 +2266,7 @@ extern "C" { GGML_API int ggml_cpu_has_gpublas (void); GGML_API int ggml_cpu_has_sse3 (void); GGML_API int ggml_cpu_has_ssse3 (void); + GGML_API int ggml_cpu_has_sycl (void); GGML_API int ggml_cpu_has_vsx (void); //