Skip to content

Commit 35cae5b

Browse files
lslusarczykRbiessy
lslusarczyk
and
Rbiessy
authored
SYCL: using graphs is configurable by environment variable and compile option (ggml-org#12371)
* alberto changes * enable sycl graphs by env variable * fixed compilation warnings in ggml-sycl.cpp * renamed graph variables * fix markdown in docs/backend/SYCL.md Co-authored-by: Romain Biessy <[email protected]> * fix markdown in docs/backend/SYCL.md again * compiling graphs by default, renamed graph_enable to graph_disable --------- Co-authored-by: Romain Biessy <[email protected]>
1 parent 810e0af commit 35cae5b

File tree

5 files changed

+56
-4
lines changed

5 files changed

+56
-4
lines changed

‎docs/backend/SYCL.md‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -660,8 +660,9 @@ use 1 SYCL GPUs: [0] with Max compute units:512
660660
|--------------------|---------------------------------------|---------------------------------------------|
661661
| GGML_SYCL | ON (mandatory) | Enable build with SYCL code path.<br>FP32 path - recommended for better perforemance than FP16 on quantized model|
662662
| GGML_SYCL_TARGET | INTEL *(default)* \| NVIDIA \| AMD | Set the SYCL target device type. |
663-
| GGML_SYCL_DEVICE_ARCH | Optional (except for AMD) | Set the SYCL device architecture, optional except for AMD. Setting the device architecture can improve the performance. See the table [--offload-arch](https://github.com/intel/llvm/blob/sycl/sycl/doc/design/OffloadDesign.md#--offload-arch) for a list of valid architectures. |
663+
| GGML_SYCL_DEVICE_ARCH | Optional (except for AMD) | Set the SYCL device architecture, optional except for AMD. Setting the device architecture can improve the performance. See the table [--offload-arch](https://github.com/intel/llvm/blob/sycl/sycl/doc/design/OffloadDesign.md#--offload-arch) for a list of valid architectures. |
664664
| GGML_SYCL_F16 | OFF *(default)* \|ON *(optional)* | Enable FP16 build with SYCL code path. |
665+
| GGML_SYCL_GRAPH | ON *(default)* \|OFF *(Optional)* | Enable build with [SYCL Graph extension](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc). |
665666
| CMAKE_C_COMPILER | `icx` *(Linux)*, `icx/cl` *(Windows)* | Set `icx` compiler for SYCL code path. |
666667
| CMAKE_CXX_COMPILER | `icpx` *(Linux)*, `icx` *(Windows)* | Set `icpx/icx` compiler for SYCL code path. |
667668

@@ -671,6 +672,7 @@ use 1 SYCL GPUs: [0] with Max compute units:512
671672
|-------------------|------------------|---------------------------------------------------------------------------------------------------------------------------|
672673
| GGML_SYCL_DEBUG | 0 (default) or 1 | Enable log function by macro: GGML_SYCL_DEBUG |
673674
| GGML_SYCL_DISABLE_OPT | 0 (default) or 1 | Disable optimize features based on Intel GPU type, to compare the performance increase |
675+
| GGML_SYCL_DISABLE_GRAPH | 0 or 1 (default) | Disable running computations through SYCL Graphs feature. Disabled by default because graph performance isn't yet better than non-graph performance. |
674676
| ZES_ENABLE_SYSMAN | 0 (default) or 1 | Support to get free memory of GPU by sycl::aspect::ext_intel_free_memory.<br>Recommended to use when --split-mode = layer |
675677

676678

‎ggml/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -186,6 +186,7 @@ option(GGML_OPENMP "ggml: use OpenMP"
186186
option(GGML_RPC "ggml: use RPC" OFF)
187187
option(GGML_SYCL "ggml: use SYCL" OFF)
188188
option(GGML_SYCL_F16 "ggml: use 16 bit floats for sycl calculations" OFF)
189+
option(GGML_SYCL_GRAPH "ggml: enable graphs in the SYCL backend" ON)
189190
set (GGML_SYCL_TARGET "INTEL" CACHE STRING
190191
"ggml: sycl target device")
191192
set (GGML_SYCL_DEVICE_ARCH "" CACHE STRING

‎ggml/src/ggml-sycl/CMakeLists.txt‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -66,6 +66,9 @@ if (WIN32)
6666
find_package(MKL REQUIRED)
6767
target_link_libraries(ggml-sycl PRIVATE IntelSYCL::SYCL_CXX MKL::MKL MKL::MKL_SYCL)
6868
else()
69+
if (GGML_SYCL_GRAPH)
70+
add_compile_definitions(GGML_SYCL_GRAPH)
71+
endif()
6972
if (GGML_SYCL_TARGET STREQUAL "INTEL")
7073
target_link_libraries(ggml-sycl PRIVATE sycl OpenCL mkl_core pthread m dl mkl_sycl_blas mkl_intel_ilp64 mkl_tbb_thread)
7174
elseif (GGML_SYCL_TARGET STREQUAL "NVIDIA")

‎ggml/src/ggml-sycl/common.hpp‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -301,6 +301,7 @@ inline optimize_feature check_gpu_optimize_feature(syclex::architecture &arch) {
301301
return opt;
302302
}
303303

304+
namespace sycl_ex = sycl::ext::oneapi::experimental;
304305
struct ggml_backend_sycl_context {
305306
int device;
306307
std::string name;
@@ -392,6 +393,10 @@ struct ggml_backend_sycl_context {
392393
return pool(device);
393394
}
394395

396+
#ifdef GGML_SYCL_GRAPH
397+
std::unique_ptr<sycl_ex::command_graph<sycl_ex::graph_state::executable>> exec_graph = nullptr;
398+
#endif
399+
395400
ggml_sycl_pool & host_pool(int device) {
396401
if (host_pools[device] == nullptr) {
397402
host_pools[device] = new_pool_for_host(stream(device, 0), device);

‎ggml/src/ggml-sycl/ggml-sycl.cpp‎

Lines changed: 44 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -46,6 +46,7 @@
4646
static bool g_sycl_loaded = false;
4747
int g_ggml_sycl_debug = 0;
4848
int g_ggml_sycl_disable_optimize = 0;
49+
int g_ggml_sycl_disable_graph = 0;
4950

5051
static ggml_sycl_device_info ggml_sycl_init() {
5152
ggml_sycl_device_info info = {};
@@ -191,10 +192,12 @@ static void ggml_check_sycl() try {
191192
if (!initialized) {
192193
g_ggml_sycl_debug = get_sycl_env("GGML_SYCL_DEBUG", 0);
193194
g_ggml_sycl_disable_optimize= get_sycl_env("GGML_SYCL_DISABLE_OPT", 0);
195+
g_ggml_sycl_disable_graph = get_sycl_env("GGML_SYCL_DISABLE_GRAPH", 1);
194196
GGML_SYCL_DEBUG("[SYCL] call ggml_check_sycl\n");
195197
GGML_LOG_INFO("Running with Environment Variables:\n");
196198
GGML_LOG_INFO(" GGML_SYCL_DEBUG: %d\n", g_ggml_sycl_debug);
197199
GGML_LOG_INFO(" GGML_SYCL_DISABLE_OPT: %d\n", g_ggml_sycl_disable_optimize);
200+
GGML_LOG_INFO(" GGML_SYCL_DISABLE_GRAPH: %d\n", g_ggml_sycl_disable_graph);
198201
GGML_LOG_INFO("Build with Macros:\n");
199202
#if defined(GGML_SYCL_FORCE_MMQ)
200203
GGML_LOG_INFO(" GGML_SYCL_FORCE_MMQ: yes\n");
@@ -3699,10 +3702,9 @@ static void optimize_graph_once(ggml_cgraph * cgraph, ggml_backend_sycl_context
36993702
if (ctx->opt_feature.reorder) opt_for_reorder(cgraph->nodes[i], stream);
37003703
}
37013704
}
3702-
static ggml_status ggml_backend_sycl_graph_compute(ggml_backend_t backend, ggml_cgraph * cgraph) {
3703-
ggml_backend_sycl_context * sycl_ctx = (ggml_backend_sycl_context *)backend->context;
3704-
ggml_sycl_set_main_device(sycl_ctx->device);
37053705

3706+
static void ggml_backend_sycl_graph_compute_impl(ggml_backend_sycl_context * sycl_ctx, ggml_cgraph * cgraph) {
3707+
ggml_sycl_set_main_device(sycl_ctx->device);
37063708
if (!g_ggml_sycl_disable_optimize) optimize_graph_once(cgraph, sycl_ctx);
37073709

37083710
for (int i = 0; i < cgraph->n_nodes; i++) {
@@ -3724,7 +3726,46 @@ static ggml_status ggml_backend_sycl_graph_compute(ggml_backend_t backend, ggml_
37243726
}
37253727
GGML_ASSERT(ok);
37263728
}
3729+
}
3730+
3731+
static ggml_status ggml_backend_sycl_graph_compute(ggml_backend_t backend, ggml_cgraph * cgraph) {
3732+
auto * sycl_ctx = static_cast<ggml_backend_sycl_context *>(backend->context);
3733+
3734+
#ifdef GGML_SYCL_GRAPH
3735+
if (!g_ggml_sycl_disable_graph) {
3736+
if (!sycl_ctx->exec_graph && !dpct::get_device(sycl_ctx->device).has(sycl::aspect::ext_oneapi_graph)) {
3737+
GGML_SYCL_DEBUG("[SYCL-GRAPH] can not use graphs on device:%d\n", sycl_ctx->device);
3738+
ggml_backend_sycl_graph_compute_impl(sycl_ctx, cgraph);
3739+
return GGML_STATUS_SUCCESS;
3740+
}
3741+
3742+
sycl_ex::command_graph model_sycl_graph(*(sycl_ctx->stream()));
3743+
model_sycl_graph.begin_recording(*(sycl_ctx->stream()));
3744+
ggml_backend_sycl_graph_compute_impl(sycl_ctx, cgraph);
3745+
model_sycl_graph.end_recording();
37273746

3747+
if (!sycl_ctx->exec_graph) {
3748+
auto exec_graph = model_sycl_graph.finalize({sycl_ex::property::graph::updatable{}});
3749+
sycl_ctx->exec_graph = std::make_unique<
3750+
sycl_ex::command_graph<sycl_ex::graph_state::executable>>(exec_graph);
3751+
} else {
3752+
try {
3753+
sycl_ctx->exec_graph->update(model_sycl_graph);
3754+
GGML_SYCL_DEBUG("[SYCL-GRAPH] update success\n");
3755+
} catch (sycl::exception const & e) {
3756+
GGML_SYCL_DEBUG("[SYCL-GRAPH] Exception when updating graph, %s\n", e.what());
3757+
auto exec_graph = model_sycl_graph.finalize({sycl_ex::property::graph::updatable{}});
3758+
sycl_ctx->exec_graph = std::make_unique<
3759+
sycl_ex::command_graph<sycl_ex::graph_state::executable>>(exec_graph);
3760+
}
3761+
}
3762+
3763+
sycl_ctx->stream()->ext_oneapi_graph(*(sycl_ctx->exec_graph));
3764+
} else
3765+
#endif
3766+
{
3767+
ggml_backend_sycl_graph_compute_impl(sycl_ctx, cgraph);
3768+
}
37283769
return GGML_STATUS_SUCCESS;
37293770
}
37303771

0 commit comments

Comments
 (0)