diff --git a/cmake/testing.cmake b/cmake/testing.cmake index a002920fcbe..f5b1314de81 100644 --- a/cmake/testing.cmake +++ b/cmake/testing.cmake @@ -32,6 +32,7 @@ set(DNNL_TEST_SET_COVERAGE "0") set(DNNL_TEST_SET_COVERAGE_STR "") set(DNNL_TEST_SET_HAS_NO_CORR "0") set(DNNL_TEST_SET_HAS_ADD_BITWISE "0") +set(DNNL_TEST_SET_HAS_SYCL_GRAPH "0") function(check_consistency entry) if(NOT DNNL_TEST_SET_COVERAGE EQUAL 0) @@ -57,6 +58,8 @@ foreach(entry ${DNNL_TEST_SET}) set(DNNL_TEST_SET_HAS_NO_CORR "1") elseif(entry STREQUAL "ADD_BITWISE") set(DNNL_TEST_SET_HAS_ADD_BITWISE "1") + elseif(entry STREQUAL "SYCL_GRAPH") + set(DNNL_TEST_SET_HAS_SYCL_GRAPH "1") elseif(entry STREQUAL "CI_NO_CORR") # Left here for compatibility till v4.0 set(DNNL_TEST_SET_COVERAGE ${DNNL_TEST_SET_CI}) set(DNNL_TEST_SET_COVERAGE_STR "CI") @@ -68,7 +71,7 @@ foreach(entry ${DNNL_TEST_SET}) message(FATAL_ERROR "The DNNL_TEST_SET entry ${entry} is not recognized. " "Supported values are:" - "NIGHTLY, CI, SMOKE, NO_CORR, ADD_BITWISE.") + "NIGHTLY, CI, SMOKE, NO_CORR, ADD_BITWISE, SYCL_GRAPH.") endif() endforeach() @@ -79,3 +82,6 @@ endif() if(DNNL_TEST_SET_HAS_ADD_BITWISE EQUAL 1) message(STATUS "Enabled testing modifier: Add bitwise validation") endif() +if(DNNL_TEST_SET_HAS_SYCL_GRAPH EQUAL 1) + message(STATUS "Enabled testing modifier: Use experimental SyCL graph execution") +endif() diff --git a/tests/benchdnn/CMakeLists.txt b/tests/benchdnn/CMakeLists.txt index 39bf83f8e4f..3f998732c51 100644 --- a/tests/benchdnn/CMakeLists.txt +++ b/tests/benchdnn/CMakeLists.txt @@ -85,6 +85,10 @@ function(register_benchdnn_test engine driver test_file) set(cmd "--mode=${tm} ${mode_modifier} -v1 --engine=${engine} --${driver} --batch=${test_file}") set(benchdnn_target ${target_name}_${engine}) + if(DNNL_TEST_SET_HAS_SYCL_GRAPH EQUAL 1) + string(PREPEND cmd "--use-sycl-graph=true") + endif() + if(NOT WIN32 OR DNNL_BUILD_FOR_CI) string(REPLACE " " ";" cmd "benchdnn ${cmd}") add_dnnl_test(${benchdnn_target} ${cmd}) diff --git a/tests/benchdnn/benchdnn.cpp b/tests/benchdnn/benchdnn.cpp index 6c32443d597..d1cb8d64b4a 100644 --- a/tests/benchdnn/benchdnn.cpp +++ b/tests/benchdnn/benchdnn.cpp @@ -56,6 +56,7 @@ int verbose {0}; bool canonical {false}; bool mem_check {true}; +bool use_sycl_graph {false}; std::string skip_impl; stat_t benchdnn_stat {0}; std::string driver_name; diff --git a/tests/benchdnn/common.hpp b/tests/benchdnn/common.hpp index 24b2c10f86d..89e14b43bba 100644 --- a/tests/benchdnn/common.hpp +++ b/tests/benchdnn/common.hpp @@ -93,6 +93,7 @@ enum { CRIT = 1, WARN = 2 }; extern int verbose; extern bool canonical; extern bool mem_check; +extern bool use_sycl_graph; extern bool attr_same_pd_check; extern bool check_ref_impl; extern std::string skip_impl; /* empty or "" means skip nothing */ diff --git a/tests/benchdnn/dnn_types.cpp b/tests/benchdnn/dnn_types.cpp index 804d133a5b2..312bf0889e0 100644 --- a/tests/benchdnn/dnn_types.cpp +++ b/tests/benchdnn/dnn_types.cpp @@ -919,6 +919,10 @@ std::ostream &dump_global_params(std::ostream &s) { #endif if (canonical || cold_cache_mode != default_cold_cache_mode) s << "--cold-cache=" << cold_cache_mode << " "; +#if defined(DNNL_WITH_SYCL) + if (canonical || use_sycl_graph != false) + s << "--use-sycl-graph=" << bool2str(use_sycl_graph) << " "; +#endif return s; } diff --git a/tests/benchdnn/dnnl_common.cpp b/tests/benchdnn/dnnl_common.cpp index 5e6de876cc5..793ec48433d 100644 --- a/tests/benchdnn/dnnl_common.cpp +++ b/tests/benchdnn/dnnl_common.cpp @@ -366,8 +366,40 @@ int execute_and_wait(perf_function_t &exec_func, const dnnl_engine_t &engine, execute_unmap_args(args, dnnl_args); - auto status = exec_func(stream, dnnl_args); - DNN_SAFE(dnnl_stream_wait(stream), CRIT); + dnnl_status_t status = dnnl_runtime_error; + bool run_regular_exec = true; +#if DNNL_GPU_RUNTIME == DNNL_RUNTIME_DPCPP + while (use_sycl_graph && is_gpu(engine)) { + void *queue_ptr; + DNN_SAFE(dnnl_sycl_interop_stream_get_queue(stream, &queue_ptr), CRIT); + sycl::queue queue = *static_cast(queue_ptr); + const bool can_run_sycl_graph = queue.get_device().get_backend() + == sycl::backend::ext_oneapi_level_zero; + if (!can_run_sycl_graph) break; + + BENCHDNN_PRINT( + 2, "%s\n", "[INFO] Using experimental SyCL graph execution."); + sycl::ext::oneapi::experimental::command_graph graph { + queue.get_context(), queue.get_device()}; + + graph.begin_recording(queue); + status = exec_func(stream, dnnl_args); + graph.end_recording(queue); + DNN_SAFE(dnnl_stream_wait(stream), CRIT); + + auto exec = graph.finalize(); + queue.ext_oneapi_graph(exec).wait(); + + // SyCL graph feature completed submission and execution, no need to + // have a regular run. + run_regular_exec = false; + break; + } +#endif + if (run_regular_exec) { + status = exec_func(stream, dnnl_args); + DNN_SAFE(dnnl_stream_wait(stream), CRIT); + } if (res) res->state = EXECUTED; execute_map_args(args); diff --git a/tests/benchdnn/utils/parser.cpp b/tests/benchdnn/utils/parser.cpp index bee2f8ab7b5..c2c29dba0b4 100644 --- a/tests/benchdnn/utils/parser.cpp +++ b/tests/benchdnn/utils/parser.cpp @@ -1378,6 +1378,26 @@ static bool parse_verbose( return false; } +static bool parse_use_sycl_graph( + const char *str, const std::string &option_name = "use-sycl-graph") { + static const std::string help + = "BOOL (Default: `false`)\n Instructs the driver to execute " + "using experimental SyCL Graph backend, when set to `true`.\n"; + bool parsed = parse_single_value_option( + use_sycl_graph, false, str2bool, str, option_name, help); + +#if !defined(DNNL_WITH_SYCL) + if (parsed) { + BENCHDNN_PRINT(0, + "Error: option `--%s` is supported with DPC++ " + "builds only, exiting...\n", + option_name.c_str()); + SAFE_V(FAIL); + } +#endif + return parsed; +} + bool parse_bench_settings(const char *str) { last_parsed_is_problem = false; // if start parsing, expect an option @@ -1402,7 +1422,8 @@ bool parse_bench_settings(const char *str) { || parse_num_streams(str) || parse_repeats_per_prb(str) || parse_mem_check(str) || parse_memory_kind(str) || parse_mode(str) || parse_mode_modifier(str) || parse_start(str) - || parse_stream_kind(str) || parse_verbose(str); + || parse_stream_kind(str) || parse_verbose(str) + || parse_use_sycl_graph(str); // Last condition makes this help message to be triggered once driver_name // is already known.