Add barrier benchmark

Multiple enqueues with barriers inbetween, wait for the event from last
barrier.

Related-To: NEO-8147

Signed-off-by: Dominik Dabek <[email protected]>

Author: ddabek-i

TESTS.md

@@ -223,6 +223,7 @@ CompletionLatency|enqueues system memory write and measures time between the mom
 CopySubmissionEvents|enqueues 4 byte copy to copy engine and return submission delta which is time between host API call and copy engine start|<ul><li>--engine Engine used for copying (RCS or CCS0 or CCS1 or CCS2 or CCS3 or BCS or BCS1 or BCS2 or BCS3 or BCS4 or BCS5 or BCS6 or BCS7 or BCS8)</li></ul>|:heavy_check_mark:|:heavy_check_mark:|
 EmptyKernel|enqueues empty kernel and measures time to launch it and wait for it on CPU, thus measuring walker spawn time.|<ul><li>--wgc Workgroup count</li><li>--wgs Workgroup size (aka local work size)</li></ul>|:heavy_check_mark:|:heavy_check_mark:|
 EmptyKernelImmediate|enqueues empty kernel and measures time to launch it using immediate command list and wait for it on CPU, thus measuring walker spawn time.|<ul><li>--UseEventForHostSync If true, use events to synchronize with host.If false, use zeCommandListHostSynchronize (0 or 1)</li><li>--wgc Workgroup count</li><li>--wgs Workgroup size (aka local work size)</li></ul>|:heavy_check_mark:|:x:|
+EnqueueBarrierWithEmptyWaitlist|enqueues kernel with barriers with empty waitlists inbetween, waiting on the last barriers event|<ul><li>--enqueueCount Number of enqueues</li><li>--outOfOrderQueue Use out of order queue (0 or 1)</li></ul>|:x:|:heavy_check_mark:|
 KernelSwitchLatency|measures time from end of one kernel till start of next kernel|<ul><li>--barrier synchronization with barrier instead of events (0 or 1)</li><li>--flush Flush between kernels (0 or 1)</li><li>--hostVisible events are with host visible flag (0 or 1)</li><li>--kernelCount Count of kernels</li><li>--kernelExecutionTime Approximately how long a single kernel executes, in us</li></ul>|:heavy_check_mark:|:heavy_check_mark:|
 KernelSwitchLatencyImmediate|measures time from end of one kernel till start of next kernel using immediate command lists|<ul><li>--barrier synchronization with barrier instead of events (0 or 1)</li><li>--hostVisible events are with host visible flag (0 or 1)</li><li>--kernelCount Count of kernels</li><li>--kernelExecutionTime Approximately how long a single kernel executes, in us</li></ul>|:heavy_check_mark:|:x:|
 KernelWithWork|measures time required to run a GPU kernel which assigns constant values to elements of a buffer. Each thread assigns one value.|<ul><li>--usedIds Which of the get_global_id() and get_local_id() calls will be used in the kernel (None or Global or Local or AtomicPerWkg)</li><li>--wgc Workgroup count</li><li>--wgs Workgroup size (aka local work size)</li></ul>|:heavy_check_mark:|:heavy_check_mark:|

source/benchmarks/ulls_benchmark/definitions/enqueue_barrier_with_empty_waitlist.h

@@ -0,0 +1,33 @@
+/*
+ * Copyright (C) 2023 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#pragma once
+
+#include "framework/argument/basic_argument.h"
+#include "framework/test_case/test_case.h"
+#include "framework/utility/common_help_message.h"
+
+struct EnqueueBarrierWithEmptyWaitlistArguments : TestCaseArgumentContainer {
+    PositiveIntegerArgument enqueueCount;
+    BooleanArgument outOfOrderQueue;
+
+    EnqueueBarrierWithEmptyWaitlistArguments()
+        : enqueueCount(*this, "enqueueCount", "Number of enqueues"),
+          outOfOrderQueue(*this, "outOfOrderQueue", "Use out of order queue") {}
+};
+
+struct EnqueueBarrierWithEmptyWaitlist : TestCase<EnqueueBarrierWithEmptyWaitlistArguments> {
+    using TestCase<EnqueueBarrierWithEmptyWaitlistArguments>::TestCase;
+
+    std::string getTestCaseName() const override {
+        return "EnqueueBarrierWithEmptyWaitlist";
+    }
+
+    std::string getHelp() const override {
+        return "enqueues kernel with barriers with empty waitlists inbetween, waiting on the last barriers event";
+    }
+};

source/benchmarks/ulls_benchmark/gtest/enqueue_barrier_with_empty_waitlist.cpp

@@ -0,0 +1,35 @@
+/*
+ * Copyright (C) 2023 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "definitions/enqueue_barrier_with_empty_waitlist.h"
+
+#include "framework/test_case/register_test_case.h"
+#include "framework/utility/common_gtest_args.h"
+
+#include <gtest/gtest.h>
+
+[[maybe_unused]] static const inline RegisterTestCase<EnqueueBarrierWithEmptyWaitlist> registerTestCase{};
+
+class EnqueueBarrierWithEmptyWaitlistTest : public ::testing::TestWithParam<std::tuple<size_t, bool>> {
+};
+
+TEST_P(EnqueueBarrierWithEmptyWaitlistTest, Test) {
+    EnqueueBarrierWithEmptyWaitlistArguments args{};
+    args.api = Api::OpenCL;
+    args.enqueueCount = std::get<0>(GetParam());
+    args.outOfOrderQueue = std::get<1>(GetParam());
+
+    EnqueueBarrierWithEmptyWaitlist test;
+    test.run(args);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    EnqueueBarrierWithEmptyWaitlistTest,
+    EnqueueBarrierWithEmptyWaitlistTest,
+    ::testing::Combine(
+        ::testing::Values(16, 32, 64),
+        ::testing::Bool()));

source/benchmarks/ulls_benchmark/implementations/ocl/enqueue_barrier_with_empty_waitlist_ocl.cpp

@@ -0,0 +1,92 @@
+/*
+ * Copyright (C) 2023 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "framework/ocl/opencl.h"
+#include "framework/ocl/utility/program_helper_ocl.h"
+#include "framework/test_case/register_test_case.h"
+#include "framework/utility/timer.h"
+
+#include "definitions/enqueue_barrier_with_empty_waitlist.h"
+
+#include <cstring>
+#include <gtest/gtest.h>
+
+static TestResult run(const EnqueueBarrierWithEmptyWaitlistArguments &arguments, Statistics &statistics) {
+    MeasurementFields typeSelector(MeasurementUnit::Microseconds, MeasurementType::Cpu);
+
+    if (isNoopRun()) {
+        statistics.pushUnitAndType(typeSelector.getUnit(), typeSelector.getType());
+        return TestResult::Nooped;
+    }
+
+    // Setup
+    QueueProperties queueProperties = QueueProperties::create().setOoq(arguments.outOfOrderQueue);
+    Opencl opencl(queueProperties);
+    Timer timer{};
+    cl_int retVal{};
+
+    // Prepare data
+    const size_t workgroupCount = 128;
+    const size_t lws = 128;
+    const size_t gws = lws * workgroupCount;
+    const size_t enqueueCount = arguments.enqueueCount;
+    if (enqueueCount == 0) {
+        return TestResult::InvalidArgs;
+    }
+
+    // Create kernel
+    cl_program program = nullptr;
+    const char *programName = "ulls_benchmark_eat_time.cl";
+    const char *kernelName = "eat_time";
+    if (auto result = ProgramHelperOcl::buildProgramFromSourceFile(opencl.context, opencl.device, programName, nullptr, program); result != TestResult::Success) {
+        return result;
+    }
+    cl_kernel kernel = clCreateKernel(program, kernelName, &retVal);
+    ASSERT_CL_SUCCESS(retVal);
+    const cl_int operationsCount = 1;
+    ASSERT_CL_SUCCESS(clSetKernelArg(kernel, 0, sizeof(size_t), &operationsCount));
+
+    // Warmup
+    {
+        cl_event event{};
+        for (auto j = 0u; j < enqueueCount - 1; ++j) {
+            ASSERT_CL_SUCCESS(clEnqueueNDRangeKernel(opencl.commandQueue, kernel, 1, nullptr, &gws, &lws, 0, nullptr, nullptr));
+            ASSERT_CL_SUCCESS(clEnqueueBarrierWithWaitList(opencl.commandQueue, 0, nullptr, nullptr));
+        }
+        ASSERT_CL_SUCCESS(clEnqueueNDRangeKernel(opencl.commandQueue, kernel, 1, nullptr, &gws, &lws, 0, nullptr, nullptr));
+        ASSERT_CL_SUCCESS(clEnqueueBarrierWithWaitList(opencl.commandQueue, 0, nullptr, &event));
+
+        ASSERT_CL_SUCCESS(clWaitForEvents(1, &event));
+        ASSERT_CL_SUCCESS(clReleaseEvent(event));
+    }
+
+    // Benchmark
+    for (auto i = 0u; i < arguments.iterations; i++) {
+        timer.measureStart();
+        {
+            cl_event event{};
+            for (auto j = 0u; j < enqueueCount - 1; ++j) {
+                ASSERT_CL_SUCCESS(clEnqueueNDRangeKernel(opencl.commandQueue, kernel, 1, nullptr, &gws, &lws, 0, nullptr, nullptr));
+                ASSERT_CL_SUCCESS(clEnqueueBarrierWithWaitList(opencl.commandQueue, 0, nullptr, nullptr));
+            }
+            ASSERT_CL_SUCCESS(clEnqueueNDRangeKernel(opencl.commandQueue, kernel, 1, nullptr, &gws, &lws, 0, nullptr, nullptr));
+            ASSERT_CL_SUCCESS(clEnqueueBarrierWithWaitList(opencl.commandQueue, 0, nullptr, &event));
+
+            ASSERT_CL_SUCCESS(clWaitForEvents(1, &event));
+            ASSERT_CL_SUCCESS(clReleaseEvent(event));
+        }
+        timer.measureEnd();
+        statistics.pushValue(timer.get(), typeSelector.getUnit(), typeSelector.getType());
+    }
+
+    // Cleanup
+    ASSERT_CL_SUCCESS(clReleaseKernel(kernel));
+    ASSERT_CL_SUCCESS(clReleaseProgram(program));
+    return TestResult::Success;
+}
+
+static RegisterTestCaseImplementation<EnqueueBarrierWithEmptyWaitlist> registerTestCase(run, Api::OpenCL);