add support for normal cmdqueue dispatch for graphs

Signed-off-by: Piotr Balcer <[email protected]>

Author: pbalcer

TESTS.md

@@ -111,6 +111,14 @@ WriteTimestamp|measures time required to write a timestamp on GPU.|<ul><li>--mea
 
 
 
+# graph_api_benchmark
+Graph Api Overhead Benchmark is a set of tests aimed at measuring CPU-side execution duration of SYCL Graphs API calls.
+| Test name | Description | Params | L0 | OCL |
+|-----------|-------------|--------|----|-----|
+SinKernelGraph|Benchmark running memory copy and kernel runs, with graphs and without graphs|<ul><li>--immediateAppendCmdList Use zeCommandListImmediateAppendCommandListsExp to submit graph (only valid for L0) (0 or 1)</li><li>--numKernels Number of kernel invocations</li><li>--withCopyOffload Enable driver copy offload (only valid for L0) (0 or 1)</li><li>--withGraphs Runs with or without graphs (0 or 1)</li></ul>|:heavy_check_mark:|:x:|
+
+
+
 # memory_benchmark
 Memory Benchmark is a set of tests aimed at measuring bandwidth of memory transfers.
 | Test name | Description | Params | L0 | OCL |

source/benchmarks/graph_api_benchmark/CMakeLists.txt

@@ -4,4 +4,4 @@
 # SPDX-License-Identifier: MIT
 #
 
-add_benchmark(graph_api_benchmark sycl all)
+add_benchmark(graph_api_benchmark sycl l0 ur all)

source/benchmarks/graph_api_benchmark/definitions/sin_kernel_graph.h

@@ -9,14 +9,23 @@
 
 #include "framework/argument/basic_argument.h"
 #include "framework/test_case/test_case.h"
+#include "framework/test_case/test_result.h"
+#include "framework/utility/timer.h"
+
+#include <math.h>
+#include <random>
 
 struct SinKernelGraphArguments : TestCaseArgumentContainer {
     PositiveIntegerArgument numKernels;
     BooleanArgument withGraphs;
+    BooleanArgument withCopyOffload;
+    BooleanArgument immediateAppendCmdList;
 
     SinKernelGraphArguments()
         : numKernels(*this, "numKernels", "Number of kernel invocations"),
-          withGraphs(*this, "withGraphs", "Runs with or without graphs") {}
+          withGraphs(*this, "withGraphs", "Runs with or without graphs"),
+          withCopyOffload(*this, "withCopyOffload", "Enable driver copy offload (only valid for L0)"),
+          immediateAppendCmdList(*this, "immediateAppendCmdList", "Use zeCommandListImmediateAppendCommandListsExp to submit graph (only valid for L0)") {}
 };
 
 struct SinKernelGraph : TestCase<SinKernelGraphArguments> {
@@ -27,6 +36,6 @@ struct SinKernelGraph : TestCase<SinKernelGraphArguments> {
     }
 
     std::string getHelp() const override {
-        return "Benchmark calling sycl::sin kernel & doing mem alloc/dealloc, with graphs and without graphs";
+        return "Benchmark running memory copy and kernel runs, with graphs and without graphs";
     }
 };

source/benchmarks/graph_api_benchmark/definitions/sin_kernel_graph_base.h

@@ -0,0 +1,170 @@
+/*
+ * Copyright (C) 2024-2025 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#pragma once
+
+#include "framework/benchmark_info.h"
+
+#include "sin_kernel_graph.h"
+
+#include <cmath>
+#include <functional>
+#include <iostream>
+#include <memory>
+#include <random>
+#include <vector>
+
+class SinKernelGraphBase {
+  public:
+    SinKernelGraphBase(const SinKernelGraphArguments &arguments)
+        : numKernels(arguments.numKernels), size(65536), withGraphs(arguments.withGraphs), withCopyOffload(arguments.withCopyOffload), immediateAppendCmdList(arguments.immediateAppendCmdList), iterations(arguments.iterations), engine(0), distribution(-10.0, 10.0){};
+
+    using DataFloatPtr = std::unique_ptr<float, std::function<void(float *)>>;
+
+    virtual DataFloatPtr allocDevice(uint32_t count) = 0;
+    virtual DataFloatPtr allocHost(uint32_t count) = 0;
+
+    virtual TestResult init() = 0;
+
+    virtual TestResult recordGraph() = 0;
+    virtual TestResult readResults(float *output_h) = 0;
+
+    virtual TestResult runGraph(float *input_h) = 0;
+    virtual TestResult runEager(float *input_h) = 0;
+    virtual TestResult waitCompletion() = 0;
+
+    TestResult calcRefResults(float *input_h, float *golden_h) {
+        std::vector<float> buffer0(size);
+        std::vector<float> buffer1(size);
+
+        // assign action
+        for (uint32_t i = 0; i < size; ++i)
+            buffer0[i] = input_h[i];
+
+        // repeat sin action
+        for (size_t k = 0; k < numKernels; ++k) {
+            std::swap(buffer0, buffer1);
+            for (uint32_t i = 0; i < size; ++i)
+                buffer0[i] = sin(buffer1[i]);
+        }
+
+        if (numKernels % 2 != 0)
+            std::swap(buffer0, buffer1);
+
+        for (uint32_t i = 0; i < size; ++i)
+            golden_h[i] = buffer0[i];
+
+        return TestResult::Success;
+    }
+
+    bool checkResults(float *output_h, float *golden_h) {
+        bool ret = true;
+        for (uint32_t idx = 0; idx < size; ++idx) {
+            if ((fabs(output_h[idx] - golden_h[idx]) > 0.00001f) ||
+                (output_h[idx]) == 0.0f) {
+                ret = false;
+                std::cout << "at (" << idx << "), expected " << golden_h[idx]
+                          << ", but got " << output_h[idx] << std::endl;
+                return ret;
+            }
+        }
+        return ret;
+    }
+
+    float randFloat() {
+        return distribution(engine);
+    }
+
+    TestResult run(
+        Statistics &statistics) {
+        MeasurementFields typeSelector(MeasurementUnit::Microseconds,
+                                       MeasurementType::Cpu);
+        if (isNoopRun()) {
+            statistics.pushUnitAndType(typeSelector.getUnit(),
+                                       typeSelector.getType());
+            return TestResult::Nooped;
+        }
+
+        init();
+
+        TestResult result = TestResult::Success;
+
+        Timer timer;
+
+        DataFloatPtr inputData = allocHost(size);
+        DataFloatPtr refResult = allocHost(size);
+        DataFloatPtr outputData = allocHost(size);
+
+        graphInputData = allocDevice(size);
+        graphOutputData = allocDevice(size);
+
+        for (uint32_t i = 0; i < size; ++i) {
+            inputData.get()[i] = randFloat();
+        }
+
+        // reference results
+        ASSERT_TEST_RESULT_SUCCESS(calcRefResults(inputData.get(), refResult.get()));
+
+        if (withGraphs)
+            ASSERT_TEST_RESULT_SUCCESS(recordGraph());
+
+        // warm-up & results verification
+        {
+            if (withGraphs) {
+                ASSERT_TEST_RESULT_SUCCESS(runGraph(inputData.get()));
+            } else {
+                ASSERT_TEST_RESULT_SUCCESS(runEager(inputData.get()));
+            }
+            ASSERT_TEST_RESULT_SUCCESS(waitCompletion());
+            ASSERT_TEST_RESULT_SUCCESS(readResults(outputData.get()));
+
+            // if results don't match, fail the benchmark
+            if (!checkResults(outputData.get(), refResult.get())) {
+                std::cout << "Check FAILED" << std::endl;
+                return TestResult::Error;
+            }
+        }
+
+        for (uint32_t i = 0; i < iterations; ++i) {
+            timer.measureStart();
+
+            if (withGraphs) {
+                ASSERT_TEST_RESULT_SUCCESS(runGraph(inputData.get()));
+            } else {
+                ASSERT_TEST_RESULT_SUCCESS(runEager(inputData.get()));
+            }
+
+            ASSERT_TEST_RESULT_SUCCESS(waitCompletion());
+
+            timer.measureEnd();
+            statistics.pushValue(timer.get(), typeSelector.getUnit(),
+                                 typeSelector.getType());
+        }
+
+        return result;
+    }
+
+    ~SinKernelGraphBase() = default;
+
+  protected:
+    size_t numKernels;
+    uint32_t size;
+    bool withGraphs;
+    bool withCopyOffload;
+    bool immediateAppendCmdList;
+
+    size_t iterations;
+
+    float pattern = 123.4567f;
+
+    // device memory
+    DataFloatPtr graphInputData;
+    DataFloatPtr graphOutputData;
+
+    std::mt19937 engine;
+    std::uniform_real_distribution<float> distribution;
+};

source/benchmarks/graph_api_benchmark/gtest/sin_kernel_graph.cpp

@@ -14,22 +14,26 @@
 
 [[maybe_unused]] static const inline RegisterTestCase<SinKernelGraph> registerTestCase{};
 
-class SinKernelGraphTest : public ::testing::TestWithParam<std::tuple<std::size_t, bool>> {
-};
+class SinKernelGraphTest
+    : public ::testing::TestWithParam<std::tuple<Api, uint32_t, bool, bool, bool>> {};
 
 TEST_P(SinKernelGraphTest, Test) {
     SinKernelGraphArguments args{};
-    args.api = Api::SYCL;
-    args.numKernels = std::get<0>(GetParam());
-    args.withGraphs = std::get<1>(GetParam());
+    args.api = std::get<0>(GetParam());
+    args.numKernels = std::get<1>(GetParam());
+    args.immediateAppendCmdList = std::get<2>(GetParam());
+    args.withCopyOffload = std::get<3>(GetParam());
+    args.withGraphs = std::get<4>(GetParam());
 
     SinKernelGraph test;
     test.run(args);
 }
 
 INSTANTIATE_TEST_SUITE_P(
-    SinKernelGraphTest,
-    SinKernelGraphTest,
-    ::testing::Combine(
-        ::testing::Values(20, 50, 100, 500),
-        ::testing::Values(false, true)));
+    SinKernelGraphTest, SinKernelGraphTest,
+    ::testing::Combine(::testing::Values(Api::SYCL, Api::UR, Api::L0),
+                       ::testing::Values(3, 10, 50, 100),
+                       // FIXME: immediateAppendCmdList is currently broken, add 'true' to enable it once the driver is fixed.
+                       ::testing::Values(false),
+                       ::testing::Values(false, true),
+                       ::testing::Values(false, true)));

source/benchmarks/graph_api_benchmark/gtest/submit_exec_graph.cpp

@@ -33,5 +33,5 @@ INSTANTIATE_TEST_SUITE_P(
     SubmitExecGraphTest,
     ::testing::Combine(
         ::testing::Values(false, true),
-        ::testing::Values(50, 100, 500),
+        ::testing::Values(20, 100),
         ::testing::Values(false, true)));

source/benchmarks/graph_api_benchmark/implementations/l0/sin_kernel_graph_l0.cpp

@@ -0,0 +1,20 @@
+/*
+ * Copyright (C) 2024-2025 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "framework/test_case/register_test_case.h"
+
+#include "definitions/sin_kernel_graph.h"
+#include "sin_kernel_impl_l0.h"
+
+static TestResult run(const SinKernelGraphArguments &arguments,
+                      Statistics &statistics) {
+    auto sin = SinKernelGraphL0(arguments);
+    return sin.run(statistics);
+}
+
+[[maybe_unused]] static RegisterTestCaseImplementation<SinKernelGraph>
+    registerTestCase(run, Api::L0);