tests: benchdnn: add experimental support for sycl graph

Author: spalicki

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()

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})

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;

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 */

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;
 }

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<sycl::queue *>(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);

tests/benchdnn/doc/knobs_common.md

@@ -146,6 +146,12 @@ Additional information is printed to the stdout depending on a level `N`. `N` is
 a non-negative integer value. The default value is `0`. Refer to
 [verbose](knobs_verbose.md) for details.
 
+### --use-sycl-graph
+`--use-sycl-graph=BOOL` instructs the driver to execute using the experimental
+SYCL Graph backend when `BOOL` is set to `true`. When `BOOL` is set to `false`
+(the default), the driver will execute normally. This feature is limited to
+DPC++ runtime with Level Zero backend.
+
 ## Correctness mode settings
 
 ### --attr-same-pd-check

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.