tests: benchdnn: add support for graph execution

Currently limited to SYCL with L0 backend.

cmake/testing.cmake

@@ -1,5 +1,5 @@
 #===============================================================================
-# Copyright 2020-2024 Intel Corporation
+# Copyright 2020-2025 Intel Corporation
 #
 # Licensed under the Apache License, Version 2.0 (the "License");
 # you may not use this file except in compliance with the License.
@@ -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_GRAPH_EXE "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 "GRAPH_EXE")
+        set(DNNL_TEST_SET_HAS_GRAPH_EXE "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, GRAPH_EXE.")
     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_GRAPH_EXE EQUAL 1)
+    message(STATUS "Enabled testing modifier: Use graph execution")
+endif()

tests/benchdnn/CMakeLists.txt

@@ -1,5 +1,5 @@
 #===============================================================================
-# Copyright 2017-2024 Intel Corporation
+# Copyright 2017-2025 Intel Corporation
 #
 # Licensed under the Apache License, Version 2.0 (the "License");
 # you may not use this file except in compliance with the License.
@@ -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_GRAPH_EXE EQUAL 1)
+            string(PREPEND cmd "--execution-mode=graph")
+        endif()
+
         if(NOT WIN32 OR DNNL_BUILD_FOR_CI)
             string(REPLACE " " ";" cmd "benchdnn ${cmd}")
             add_dnnl_test(${benchdnn_target} ${cmd})

tests/benchdnn/benchdnn.cpp

@@ -1,5 +1,5 @@
 /*******************************************************************************
-* Copyright 2017-2024 Intel Corporation
+* Copyright 2017-2025 Intel Corporation
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
@@ -76,6 +76,8 @@ int test_start {0};
 bool attr_same_pd_check {false};
 bool check_ref_impl {false};
 
+execution_mode_t execution_mode {execution_mode_t::direct};
+
 int main(int argc, char **argv) {
     using namespace parser;
 

tests/benchdnn/dnn_types.cpp

@@ -920,6 +920,8 @@ 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 (canonical || execution_mode != execution_mode_t::direct)
+        s << "--execution-mode=" << execution_mode2str(execution_mode) << " ";
 
     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);
+    dnnl_status_t status = dnnl_runtime_error;
+    bool run_regular_exec = true;
+#if DNNL_GPU_RUNTIME == DNNL_RUNTIME_DPCPP
+    while (execution_mode == execution_mode_t::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);
@@ -1371,6 +1403,32 @@ memory_kind_ext_t str2memory_kind(const char *str) {
     return memory_kind_ext_t::usm;
 }
 
+const char *execution_mode2str(execution_mode_t mode) {
+#define EXECUTION_MODE_TO_STR(name, ...) \
+    if (execution_mode_t::name == mode) return #name;
+
+    EXECUTION_MODE_TO_STR(direct);
+    EXECUTION_MODE_TO_STR(graph);
+#undef EXECUTION_MODE_STR
+
+    BENCHDNN_PRINT(0, "%s", "Error: execution mode value is not recognized.\n");
+    SAFE_V(FAIL);
+    return "";
+}
+
+execution_mode_t str2execution_mode(const char *str) {
+#define STR_TO_EXECUTION_MODE(name, ...) \
+    if (!strcasecmp(#name, str)) return execution_mode_t::name;
+
+    STR_TO_EXECUTION_MODE(direct);
+    STR_TO_EXECUTION_MODE(graph);
+#undef STR_TO_EXECUTION_MODE
+
+    BENCHDNN_PRINT(0, "%s", "Error: execution mode value is not recognized.\n");
+    SAFE_V(FAIL);
+    return execution_mode_t::direct;
+}
+
 static void maybe_print_cpu_engine_error_message() {
 #if DNNL_CPU_RUNTIME == DNNL_RUNTIME_SYCL
     fprintf(stderr,

tests/benchdnn/dnnl_common.hpp

@@ -1,5 +1,5 @@
 /*******************************************************************************
-* Copyright 2017-2024 Intel Corporation
+* Copyright 2017-2025 Intel Corporation
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
@@ -655,6 +655,12 @@ bool check_md_consistency_with_tag(
 
 memory_kind_ext_t str2memory_kind(const char *str);
 
+enum class execution_mode_t { direct, graph };
+extern execution_mode_t execution_mode;
+
+const char *execution_mode2str(execution_mode_t mode);
+execution_mode_t str2execution_mode(const char *str);
+
 float reorder_rescale_factor();
 
 // The function converts a memory descriptor dims into a `dims_t` object under

tests/benchdnn/doc/knobs_common.md

@@ -146,6 +146,13 @@ 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.
 
+### --execution-mode
+`--execution-mode=MODE` specifies the execution mode to be used. When `MODE`
+is set to `direct` (the default), the driver will execute normally. When `MODE`
+is set to `graph` it instructs the driver to execute on a graph backend.
+Currently this feature is limited to the experimental SYCL Graph feature on
+DPC++ runtime with Level Zero backend.
+
 ## Correctness mode settings
 
 ### --attr-same-pd-check

tests/benchdnn/utils/parser.cpp

@@ -1353,6 +1353,33 @@ static bool parse_verbose(
     return false;
 }
 
+static bool parse_execution_mode(
+        const char *str, const std::string &option_name = "execution-mode") {
+    static const std::string help
+            = "MODE    (Default: direct)\n"
+              "    Specifies a `MODE` of execution.\n"
+              "    `MODE` values are:\n"
+              "    * `direct` instruction the driver to execute the primitive "
+              "directly.\n"
+              "    * `graph` to execute the primitive using a graph backend.\n"
+              "          Currently limited to the experimental SYCL Graph on "
+              "DPC++ builds.\n";
+    bool parsed = parse_single_value_option(execution_mode,
+            execution_mode_t::direct, str2execution_mode, 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
 
@@ -1377,7 +1404,8 @@ bool parse_bench_settings(const char *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_execution_mode(str);
 
     // Last condition makes this help message to be triggered once driver_name
     // is already known.