Refactor Zello Metrics

1. Refactor Zello Metrics to seperate Metric Collection,
Workload Execution and Execution environment configuration
functions.
2. Split Zello metrics test to smaller files for better
readability.

Related-To: LOCI-3069

Signed-off-by: Joshua Santosh Ranjan <joshua.santosh.ranjan@intel.com>

level_zero/tools/test/black_box_tests/CMakeLists.txt

@@ -1,5 +1,5 @@
 #
-# Copyright (C) 2020-2021 Intel Corporation
+# Copyright (C) 2020-2022 Intel Corporation
 #
 # SPDX-License-Identifier: MIT
 #
@@ -8,7 +8,6 @@ if(UNIX)
   set(L0_BLACK_BOX_TEST_PROJECT_FOLDER "ze_intel_gpu/black_box_tests")
   set(TEST_TARGETS
       zello_sysman
-      zello_metrics
   )
 
   foreach(TEST_NAME ${TEST_TARGETS})

level_zero/tools/test/black_box_tests/zello_metrics/CMakeLists.txt

@@ -0,0 +1,29 @@
+#
+# Copyright (C) 2022 Intel Corporation
+#
+# SPDX-License-Identifier: MIT
+#
+
+if(UNIX)
+  set(L0_ZELLO_METRICS_PROJECT_FOLDER "ze_intel_gpu/black_box_tests")
+  set(TEST_TARGET zello_metrics)
+
+  add_executable(zello_metrics zello_metrics_collector.cpp
+                 zello_metrics_execution_context.cpp
+                 zello_metrics_workload.cpp
+                 zello_metrics.cpp
+                 zello_metrics_util.cpp
+                 zello_metrics_util.h
+                 zello_metrics.h
+  )
+
+  if(BUILD_LEVEL_ZERO_LOADER)
+    add_dependencies(${TEST_TARGET} ze_loader)
+    target_link_libraries(${TEST_TARGET} ${NEO_BINARY_DIR}/lib/libze_loader.so)
+  else()
+    target_link_libraries(${TEST_TARGET} PUBLIC ${TARGET_NAME_L0})
+  endif()
+  set_target_properties(${TEST_TARGET} PROPERTIES FOLDER ${L0_ZELLO_METRICS_PROJECT_FOLDER})
+endif()
+
+add_subdirectories()

level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics.cpp

@@ -0,0 +1,551 @@
+/*
+ * Copyright (C) 2022 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics.h"
+
+#include "level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics_util.h"
+
+#include <algorithm>
+#include <functional>
+#include <map>
+#include <memory>
+#include <mutex>
+#include <sys/wait.h>
+#include <thread>
+#include <unistd.h>
+
+////////////////////////
+/////SingleDeviceTestRunner
+////////////////////////
+void SingleDeviceTestRunner::addCollector(Collector *collector) {
+    // Add only if not available already
+    if (std::find(collectorList.begin(), collectorList.end(), collector) == collectorList.end()) {
+        collectorList.push_back(collector);
+    }
+}
+
+void SingleDeviceTestRunner::addWorkload(Workload *workload) {
+    // Add only if not available already
+    if (std::find(workloadList.begin(), workloadList.end(), workload) == workloadList.end()) {
+        workloadList.push_back(workload);
+    }
+}
+
+bool SingleDeviceTestRunner::run() {
+
+    bool status = true;
+
+    status &= executionCtxt->activateMetricGroups();
+
+    EXPECT(status == true);
+    // Command List prepareation
+    for (auto collector : collectorList) {
+        status &= collector->start();
+    }
+
+    EXPECT(status == true);
+    for (auto collector : collectorList) {
+        status &= collector->prefixCommands();
+    }
+
+    EXPECT(status == true);
+    for (auto workload : workloadList) {
+        status &= workload->appendCommands();
+    }
+
+    EXPECT(status == true);
+    for (auto collector : collectorList) {
+        status &= collector->suffixCommands();
+    }
+
+    if (disableCommandListExecution == false) {
+        EXPECT(status == true);
+        // Command List execution
+        executionCtxt->run();
+    }
+
+    EXPECT(status == true);
+    for (auto collector : collectorList) {
+        status &= collector->isDataAvailable();
+    }
+
+    EXPECT(status == true);
+    for (auto workload : workloadList) {
+        status &= workload->validate();
+    }
+
+    EXPECT(status == true);
+    for (auto collector : collectorList) {
+        collector->showResults();
+    }
+
+    for (auto collector : collectorList) {
+        status &= collector->stop();
+    }
+
+    EXPECT(status == true);
+    status &= executionCtxt->deactivateMetricGroups();
+
+    return status;
+}
+
+///////////////////////////
+/// query_device_0_sub_device_0
+///////////////////////////
+bool query_device_0_sub_device_0(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric query: device 0 ====-" << std::endl;
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(0, 0);
+    std::unique_ptr<CopyBufferToBuffer> workload =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricQueryCollector> collector =
+        std::make_unique<SingleMetricQueryCollector>(executionCtxt.get(), "TestOa");
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload.get());
+    return testRunner->run();
+}
+
+///////////////////////////
+/// query_device_root
+///////////////////////////
+bool query_device_root(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric query: device 0 ====-" << std::endl;
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(0, -1);
+    std::unique_ptr<CopyBufferToBuffer> workload =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricQueryCollector> collector =
+        std::make_unique<SingleMetricQueryCollector>(executionCtxt.get(), "TestOa");
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload.get());
+    return testRunner->run();
+}
+
+///////////////////////////
+/// query_device_0_sub_device_1
+///////////////////////////
+bool query_device_0_sub_device_1(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric query: device 0 / sub_device 1 ====-" << std::endl;
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(0, 1);
+    std::unique_ptr<CopyBufferToBuffer> workload =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricQueryCollector> collector =
+        std::make_unique<SingleMetricQueryCollector>(executionCtxt.get(), "TestOa");
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload.get());
+    return testRunner->run();
+}
+
+///////////////////////////
+/// query_device_0_sub_device_0_1
+///////////////////////////
+bool query_device_0_sub_device_0_1(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric query: device 0 / sub_device 0 : 1 ====-" << std::endl;
+
+    ZelloMetricsUtility::TestMachineConfiguration machineConfig = {};
+    ZelloMetricsUtility::getTestMachineConfiguration(machineConfig);
+
+    if (machineConfig.devices[0].subDeviceCount < 2) {
+        std::cout << "Error: Device has insufficient sub-devices" << std::endl;
+        return false;
+    }
+
+    bool status = true;
+    const char *metricGroupNames[] = {
+        "TestOa",
+        "ComputeBasic"};
+    for (int32_t subDeviceId = 0; subDeviceId < 2; subDeviceId++) {
+        std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+            std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(0, subDeviceId);
+        std::unique_ptr<CopyBufferToBuffer> workload =
+            std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+        std::unique_ptr<SingleMetricQueryCollector> collector =
+            std::make_unique<SingleMetricQueryCollector>(executionCtxt.get(), metricGroupNames[subDeviceId]);
+
+        std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+        testRunner->addCollector(collector.get());
+        testRunner->addWorkload(workload.get());
+        status &= testRunner->run();
+    }
+
+    return status;
+}
+
+///////////////////////////
+/// query_device_1_sub_device_1
+///////////////////////////
+bool query_device_1_sub_device_1(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric query: device 1 / sub_device 1 ====-" << std::endl;
+
+    ZelloMetricsUtility::TestMachineConfiguration machineConfig = {};
+    ZelloMetricsUtility::getTestMachineConfiguration(machineConfig);
+
+    if (machineConfig.deviceCount < 2) {
+        std::cout << "Error: Device has insufficient devices" << std::endl;
+        return false;
+    }
+
+    if (machineConfig.devices[0].subDeviceCount < 2) {
+        std::cout << "Error: Device has insufficient sub-devices" << std::endl;
+        return false;
+    }
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(1, 1);
+    std::unique_ptr<CopyBufferToBuffer> workload =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricQueryCollector> collector =
+        std::make_unique<SingleMetricQueryCollector>(executionCtxt.get(), "TestOa");
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload.get());
+
+    return testRunner->run();
+}
+
+///////////////////////////
+/// query_device_1_sub_device_0
+///////////////////////////
+bool query_device_1_sub_device_0(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric query: device 1 ====-" << std::endl;
+
+    ZelloMetricsUtility::TestMachineConfiguration machineConfig = {};
+    ZelloMetricsUtility::getTestMachineConfiguration(machineConfig);
+
+    if (machineConfig.deviceCount < 2) {
+        std::cout << "Error: Device has insufficient devices" << std::endl;
+        return false;
+    }
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(1, 0);
+    std::unique_ptr<CopyBufferToBuffer> workload =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricQueryCollector> collector =
+        std::make_unique<SingleMetricQueryCollector>(executionCtxt.get(), "TestOa");
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload.get());
+
+    return testRunner->run();
+}
+
+//////////////////////////////////
+/// query_device_0_1_sub_device_0
+//////////////////////////////////
+bool query_device_0_1_sub_device_0(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric query: device 0 / 1 ====-" << std::endl;
+
+    ZelloMetricsUtility::TestMachineConfiguration machineConfig = {};
+    ZelloMetricsUtility::getTestMachineConfiguration(machineConfig);
+
+    if (machineConfig.deviceCount < 2) {
+        std::cout << "Error: Device has insufficient devices" << std::endl;
+        return false;
+    }
+
+    bool status = true;
+    const char *metricGroupNames[] = {
+        "TestOa",
+        "ComputeBasic"};
+    for (int32_t deviceId = 0; deviceId < 2; deviceId++) {
+        std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+            std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(deviceId, 0);
+        std::unique_ptr<CopyBufferToBuffer> workload =
+            std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+        std::unique_ptr<SingleMetricQueryCollector> collector =
+            std::make_unique<SingleMetricQueryCollector>(executionCtxt.get(), metricGroupNames[deviceId]);
+
+        std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+        testRunner->addCollector(collector.get());
+        testRunner->addWorkload(workload.get());
+        status &= testRunner->run();
+    }
+
+    return status;
+}
+
+bool stream_device_root(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric stream: device 0 ====-" << std::endl;
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(0, -1);
+    std::unique_ptr<CopyBufferToBuffer> workload =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricStreamerCollector> collector =
+        std::make_unique<SingleMetricStreamerCollector>(executionCtxt.get(), "ComputeBasic");
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload.get());
+
+    return testRunner->run();
+    ;
+}
+
+////////////////////////////////
+/// stream_device_0_sub_device_0
+////////////////////////////////
+bool stream_device_0_sub_device_0(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric stream: device 0 ====-" << std::endl;
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(0, 0);
+    std::unique_ptr<CopyBufferToBuffer> workload =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricStreamerCollector> collector =
+        std::make_unique<SingleMetricStreamerCollector>(executionCtxt.get(), "TestOa");
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload.get());
+
+    return testRunner->run();
+}
+
+///////////////////////////
+/// stream_device_0_sub_device_1
+///////////////////////////
+bool stream_device_0_sub_device_1(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric stream: device 0 / sub_device 1 ====-" << std::endl;
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(0, 1);
+    std::unique_ptr<CopyBufferToBuffer> workload =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricStreamerCollector> collector =
+        std::make_unique<SingleMetricStreamerCollector>(executionCtxt.get(), "TestOa");
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload.get());
+
+    return testRunner->run();
+}
+
+///////////////////////////
+/// stream_device_0_sub_device_0_1
+///////////////////////////
+bool stream_device_0_sub_device_0_1(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric stream: device 0 / sub_device 0 : 1 ====-" << std::endl;
+
+    ZelloMetricsUtility::TestMachineConfiguration machineConfig = {};
+    ZelloMetricsUtility::getTestMachineConfiguration(machineConfig);
+
+    if (machineConfig.devices[0].subDeviceCount < 2) {
+        std::cout << "Error: Device has insufficient sub-devices" << std::endl;
+        return false;
+    }
+
+    bool status = true;
+    const char *metricGroupNames[] = {
+        "TestOa",
+        "ComputeBasic"};
+    for (int32_t subDeviceId = 0; subDeviceId < 2; subDeviceId++) {
+        std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+            std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(0, subDeviceId);
+        std::unique_ptr<CopyBufferToBuffer> workload =
+            std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+        std::unique_ptr<SingleMetricStreamerCollector> collector =
+            std::make_unique<SingleMetricStreamerCollector>(executionCtxt.get(), metricGroupNames[subDeviceId]);
+
+        std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+        testRunner->addCollector(collector.get());
+        testRunner->addWorkload(workload.get());
+        status &= testRunner->run();
+    }
+
+    return status;
+}
+
+/////////////////////////////////
+/// stream_device_1_sub_device_0
+/////////////////////////////////
+bool stream_device_1_sub_device_0(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric stream: device 1 ====-" << std::endl;
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(1, 0);
+    std::unique_ptr<CopyBufferToBuffer> workload =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricStreamerCollector> collector =
+        std::make_unique<SingleMetricStreamerCollector>(executionCtxt.get(), "ComputeBasic");
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload.get());
+
+    return testRunner->run();
+}
+
+///////////////////////////
+/// stream_device_1_sub_device_1
+///////////////////////////
+bool stream_device_1_sub_device_1(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric stream: device 1 / sub_device 1 ====-" << std::endl;
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(1, 1);
+    std::unique_ptr<CopyBufferToBuffer> workload =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricStreamerCollector> collector =
+        std::make_unique<SingleMetricStreamerCollector>(executionCtxt.get(), "ComputeBasic");
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload.get());
+
+    return testRunner->run();
+}
+
+///////////////////////////////////
+/// stream_device_0_1_sub_device_0
+///////////////////////////////////
+bool stream_device_0_1_sub_device_0(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric stream: device 0 / 1 ====-" << std::endl;
+
+    ZelloMetricsUtility::TestMachineConfiguration machineConfig = {};
+    ZelloMetricsUtility::getTestMachineConfiguration(machineConfig);
+
+    if (machineConfig.deviceCount < 2) {
+        std::cout << "Error: Device has insufficient devices" << std::endl;
+        return false;
+    }
+
+    bool status = true;
+    const char *metricGroupNames[] = {
+        "TestOa",
+        "ComputeBasic"};
+    for (int32_t deviceId = 0; deviceId < 2; deviceId++) {
+        std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+            std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(deviceId, 0);
+        std::unique_ptr<CopyBufferToBuffer> workload =
+            std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+        std::unique_ptr<SingleMetricStreamerCollector> collector =
+            std::make_unique<SingleMetricStreamerCollector>(executionCtxt.get(), metricGroupNames[deviceId]);
+
+        std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+        testRunner->addCollector(collector.get());
+        testRunner->addWorkload(workload.get());
+        status &= testRunner->run();
+    }
+
+    return status;
+}
+
+/////////////////////////////////////////////
+/// stream_ip_sampling_device_0_sub_device_0
+/////////////////////////////////////////////
+bool stream_ip_sampling_device_0_sub_device_0(int argc, char *argv[]) {
+
+    std::cout << std::endl
+              << "-==== Metric stream Ip Sampling: device 0 / sub_device 0 ====-" << std::endl;
+
+    std::unique_ptr<SingleDeviceSingleQueueExecutionCtxt> executionCtxt =
+        std::make_unique<SingleDeviceSingleQueueExecutionCtxt>(0, 0);
+    executionCtxt->setExecutionTimeInMilliseconds(200);
+    std::unique_ptr<AppendMemoryCopyFromHeapToDeviceAndBackToHost> workload1 =
+        std::make_unique<AppendMemoryCopyFromHeapToDeviceAndBackToHost>(executionCtxt.get());
+    std::unique_ptr<CopyBufferToBuffer> workload2 =
+        std::make_unique<CopyBufferToBuffer>(executionCtxt.get());
+    std::unique_ptr<SingleMetricStreamerCollector> collector =
+        std::make_unique<SingleMetricStreamerCollector>(executionCtxt.get(), "EuStallSampling");
+    collector->setMaxRequestRawReportCount(1000);
+
+    std::unique_ptr<SingleDeviceTestRunner> testRunner = std::make_unique<SingleDeviceTestRunner>(static_cast<ExecutionContext *>(executionCtxt.get()));
+    testRunner->addCollector(collector.get());
+    testRunner->addWorkload(workload1.get());
+    testRunner->addWorkload(workload2.get());
+
+    return testRunner->run();
+}
+
+int main(int argc, char *argv[]) {
+
+    printf("Zello metrics\n");
+    fflush(stdout);
+
+    std::map<std::string, std::function<bool(int argc, char *argv[])>> tests;
+
+    tests["query_device_root"] = query_device_root;
+    tests["query_device_0_sub_device_0"] = query_device_0_sub_device_0;
+    tests["query_device_0_sub_device_1"] = query_device_0_sub_device_1;
+    tests["query_device_0_sub_device_0_1"] = query_device_0_sub_device_0_1;
+    tests["query_device_1_sub_device_0"] = query_device_1_sub_device_0;
+    tests["query_device_1_sub_device_1"] = query_device_1_sub_device_1;
+    tests["query_device_0_1_sub_device_0"] = query_device_0_1_sub_device_0;
+
+    tests["stream_device_root"] = stream_device_root;
+    tests["stream_device_0_sub_device_0"] = stream_device_0_sub_device_0;
+    tests["stream_device_0_sub_device_1"] = stream_device_0_sub_device_1;
+    tests["stream_device_0_sub_device_0_1"] = stream_device_0_sub_device_0_1;
+    tests["stream_device_1_sub_device_0"] = stream_device_1_sub_device_0;
+    tests["stream_device_1_sub_device_1"] = stream_device_1_sub_device_1;
+    tests["stream_device_0_1_sub_device_0"] = stream_device_0_1_sub_device_0;
+    tests["stream_ip_sampling_device_0_sub_device_0"] = stream_ip_sampling_device_0_sub_device_0;
+
+    // Run test.
+    if (argc > 1) {
+        if (tests.find(argv[1]) != tests.end()) {
+            if (tests[argv[1]](argc, argv) == true) {
+                std::cout << argv[1] << " : PASS"
+                          << "\n";
+            } else {
+                std::cout << argv[1] << " : FAIL"
+                          << "\n";
+            }
+            return 0;
+        }
+    }
+
+    // Print available tests.
+    for (auto &test : tests) {
+        std::cout << test.first.c_str() << std::endl;
+    }
+
+    return 0;
+}

level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics.h

@@ -0,0 +1,250 @@
+/*
+ * Copyright (C) 2022 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include <level_zero/ze_api.h>
+#include <level_zero/zet_api.h>
+
+#include <iostream>
+#include <vector>
+
+#pragma once
+
+#define VALIDATECALL(myZeCall)                  \
+    do {                                        \
+        if ((myZeCall) != ZE_RESULT_SUCCESS) {  \
+            std::cout << "Validate Error at "   \
+                      << #myZeCall << ": "      \
+                      << __FILE__ << ": "       \
+                      << __LINE__ << std::endl; \
+            std::terminate();                   \
+        }                                       \
+    } while (0);
+
+#define EXPECT(cond)                                                           \
+    do {                                                                       \
+        if (!(cond)) {                                                         \
+            std::cout << "Error @ " << __LINE__ << " of " << __FILE__ << "\n"; \
+            std::terminate();                                                  \
+        }                                                                      \
+    } while (0);
+
+class ExecutionContext {
+
+  public:
+    ExecutionContext() = default;
+    virtual ~ExecutionContext() = default;
+    virtual bool run() = 0;
+
+    virtual ze_driver_handle_t getDriverHandle(uint32_t index) = 0;
+    virtual ze_context_handle_t getContextHandle(uint32_t index) = 0;
+    virtual ze_device_handle_t getDeviceHandle(uint32_t index) = 0;
+    virtual ze_command_queue_handle_t getCommandQueue(uint32_t index) = 0;
+    virtual ze_command_list_handle_t getCommandList(uint32_t index) = 0;
+
+    virtual void addActiveMetricGroup(zet_metric_group_handle_t &metricGroup);
+    virtual void removeActiveMetricGroup(zet_metric_group_handle_t metricGroup);
+    virtual bool activateMetricGroups();
+    virtual bool deactivateMetricGroups();
+
+  protected:
+    std::vector<zet_metric_group_handle_t> activeMetricGroups = {};
+};
+
+class Workload {
+  public:
+    Workload(ExecutionContext *execCtxt) { executionCtxt = execCtxt; }
+    virtual ~Workload() = default;
+    virtual bool appendCommands() = 0;
+    virtual bool validate() = 0;
+
+  protected:
+    ExecutionContext *executionCtxt = nullptr;
+};
+
+class Collector {
+  public:
+    Collector(ExecutionContext *execCtxt) { executionCtxt = execCtxt; }
+    virtual ~Collector() = default;
+    virtual bool prefixCommands() = 0;
+    virtual bool suffixCommands() = 0;
+    virtual bool isDataAvailable() = 0;
+    virtual bool start() = 0;
+    virtual bool stop() = 0;
+    virtual void showResults() = 0;
+
+  protected:
+    ExecutionContext *executionCtxt = nullptr;
+};
+
+class SingleDeviceSingleQueueExecutionCtxt : public ExecutionContext {
+
+  public:
+    SingleDeviceSingleQueueExecutionCtxt(uint32_t deviceIndex, int32_t subDeviceIndex = -1) { initialize(deviceIndex, subDeviceIndex); }
+    virtual ~SingleDeviceSingleQueueExecutionCtxt() { finalize(); }
+    bool run() override;
+
+    ze_driver_handle_t getDriverHandle(uint32_t index) override { return driverHandle; }
+    ze_context_handle_t getContextHandle(uint32_t index) override { return contextHandle; }
+    ze_device_handle_t getDeviceHandle(uint32_t index) override { return deviceHandle; }
+    ze_command_queue_handle_t getCommandQueue(uint32_t index) override { return commandQueue; }
+    ze_command_list_handle_t getCommandList(uint32_t index) override { return commandList; }
+
+    void setExecutionTimeInMilliseconds(uint32_t value) { executionTimeInMilliSeconds = value; }
+
+  private:
+    ze_driver_handle_t driverHandle = {};
+    ze_context_handle_t contextHandle = {};
+    ze_device_handle_t deviceHandle = {};
+    ze_command_queue_handle_t commandQueue = {};
+    ze_command_list_handle_t commandList = {};
+    bool initialize(uint32_t deviceIndex, int32_t subDeviceIndex);
+    bool finalize();
+
+    uint32_t executionTimeInMilliSeconds = 0;
+};
+
+class AppendMemoryCopyFromHeapToDeviceAndBackToHost : public Workload {
+  public:
+    AppendMemoryCopyFromHeapToDeviceAndBackToHost(ExecutionContext *execCtxt) : Workload(execCtxt) { initialize(); }
+    virtual ~AppendMemoryCopyFromHeapToDeviceAndBackToHost() { finalize(); };
+    bool appendCommands() override;
+    bool validate() override;
+
+  private:
+    void initialize();
+    void finalize();
+
+  private:
+    static constexpr size_t allocSize = 4096 + 7; // +7 to break alignment and make it harder
+    char *heapBuffer1 = nullptr;
+    char *heapBuffer2 = nullptr;
+    void *zeBuffer = nullptr;
+};
+
+class CopyBufferToBuffer : public Workload {
+  public:
+    CopyBufferToBuffer(ExecutionContext *execCtxt) : Workload(execCtxt) { initialize(); }
+    virtual ~CopyBufferToBuffer() { finalize(); };
+    bool appendCommands() override;
+    bool validate() override;
+
+  private:
+    void initialize();
+    void finalize();
+
+  private:
+    const uint32_t allocationSize = 4096;
+    void *sourceBuffer = nullptr;
+    void *destinationBuffer = nullptr;
+    ze_module_handle_t module = nullptr;
+    ze_kernel_handle_t kernel = nullptr;
+    bool executeFromSpirv = false;
+};
+
+class SingleMetricCollector : public Collector {
+
+  public:
+    void setVerboseLevel(uint32_t level) { verboseLevel = level; }
+
+  protected:
+    SingleMetricCollector(ExecutionContext *executionCtxt,
+                          const char *metricGroupName,
+                          const zet_metric_group_sampling_type_flag_t samplingType);
+    virtual ~SingleMetricCollector() = default;
+
+    virtual bool prefixCommands() override = 0;
+    virtual bool suffixCommands() override = 0;
+    virtual bool isDataAvailable() override = 0;
+    virtual bool start() override = 0;
+    virtual bool stop() override = 0;
+    virtual void showResults() override = 0;
+
+    zet_metric_group_handle_t metricGroup = {};
+    zet_metric_group_sampling_type_flag_t samplingType = {};
+
+    uint32_t verboseLevel = 0;
+};
+
+class SingleMetricStreamerCollector : public SingleMetricCollector {
+
+  public:
+    SingleMetricStreamerCollector(ExecutionContext *executionCtxt,
+                                  const char *metricGroupName);
+
+    virtual ~SingleMetricStreamerCollector() = default;
+
+    bool prefixCommands() override;
+    bool suffixCommands() override;
+    bool isDataAvailable() override;
+
+    bool start() override;
+    bool stop() override;
+
+    void setNotifyReportCount(uint32_t reportCount) { notifyReportCount = reportCount; }
+    void setSamplingPeriod(uint32_t period) { samplingPeriod = period; }
+    void setMaxRequestRawReportCount(uint32_t reportCount) { maxRequestRawReportCount = reportCount; }
+
+    virtual void showResults() override;
+
+  protected:
+    zet_metric_streamer_handle_t metricStreamer = {};
+    ze_event_pool_handle_t eventPool = {};
+    ze_event_handle_t notificationEvent = {};
+    uint32_t notifyReportCount = 1;
+    uint32_t samplingPeriod = 10000;
+    uint32_t maxRequestRawReportCount = 5;
+};
+
+class SingleMetricQueryCollector : public SingleMetricCollector {
+
+  public:
+    SingleMetricQueryCollector(ExecutionContext *executionCtxt,
+                               const char *metricGroupName);
+
+    virtual ~SingleMetricQueryCollector() = default;
+
+    bool prefixCommands() override;
+    bool suffixCommands() override;
+    bool isDataAvailable() override;
+
+    bool start() override;
+    bool stop() override;
+
+    void setAttributes(uint32_t notifyReportCount, zet_metric_query_pool_type_t queryPoolType);
+    void showResults() override;
+
+  private:
+    zet_metric_query_pool_type_t queryPoolType = ZET_METRIC_QUERY_POOL_TYPE_PERFORMANCE;
+    uint32_t queryPoolCount = 1000;
+    uint32_t querySlotIndex = 0;
+    ze_event_pool_handle_t eventPool = {};
+    ze_event_handle_t notificationEvent = {};
+    uint32_t notifyReportCount = 1;
+
+    zet_metric_query_pool_handle_t queryPoolHandle = {};
+    zet_metric_query_handle_t queryHandle = {};
+    zet_metric_query_pool_desc_t queryPoolDesc = {};
+};
+
+class SingleDeviceTestRunner {
+
+  public:
+    SingleDeviceTestRunner(ExecutionContext *context) : executionCtxt(context) {}
+    virtual ~SingleDeviceTestRunner() = default;
+
+    void addCollector(Collector *collector);
+    void addWorkload(Workload *workload);
+    // In case of streamer collector only cases
+    void disableCommandsExecution() { disableCommandListExecution = true; }
+    virtual bool run();
+
+  private:
+    std::vector<Collector *> collectorList = {};
+    std::vector<Workload *> workloadList = {};
+    ExecutionContext *executionCtxt = nullptr;
+    bool disableCommandListExecution = false;
+};

level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics_collector.cpp

@@ -0,0 +1,190 @@
+/*
+ * Copyright (C) 2022 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics.h"
+#include "level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics_util.h"
+
+#include <cstring>
+#include <thread>
+#include <vector>
+
+///////////////////////////
+/////SingleMetricCollector
+///////////////////////////
+SingleMetricCollector::SingleMetricCollector(ExecutionContext *executionCtxt,
+                                             const char *metricGroupName,
+                                             const zet_metric_group_sampling_type_flag_t samplingType)
+    : Collector(executionCtxt), samplingType(samplingType) {
+
+    metricGroup = ZelloMetricsUtility::findMetricGroup(metricGroupName, samplingType, executionCtxt->getDeviceHandle(0));
+
+    if (verboseLevel >= 1) {
+        zet_metric_group_properties_t metricGroupProperties = {};
+        // Obtain metric group properties to check the group name and sampling type.
+        VALIDATECALL(zetMetricGroupGetProperties(metricGroup, &metricGroupProperties));
+        ZelloMetricsUtility::printMetricGroupProperties(metricGroupProperties);
+
+        // Print metrics from metric group.
+        uint32_t metricCount = 0;
+        std::vector<zet_metric_handle_t> metrics = {};
+
+        // Obtain metrics count for verbose purpose.
+        VALIDATECALL(zetMetricGet(metricGroup, &metricCount, nullptr));
+
+        // Obtain metrics for verbose purpose.
+        metrics.resize(metricCount);
+        VALIDATECALL(zetMetricGet(metricGroup, &metricCount, metrics.data()));
+
+        // Enumerate metric group metrics for verbose purpose.
+        for (uint32_t j = 0; j < metricCount; ++j) {
+
+            const zet_metric_handle_t metric = metrics[j];
+            zet_metric_properties_t metricProperties = {};
+            VALIDATECALL(zetMetricGetProperties(metric, &metricProperties));
+            ZelloMetricsUtility::printMetricProperties(metricProperties);
+        }
+    }
+}
+
+///////////////////////////////////
+/////SingleMetricStreamerCollector
+///////////////////////////////////
+SingleMetricStreamerCollector::SingleMetricStreamerCollector(ExecutionContext *executionCtxt,
+                                                             const char *metricGroupName) : SingleMetricCollector(executionCtxt, metricGroupName,
+                                                                                                                  ZET_METRIC_GROUP_SAMPLING_TYPE_FLAG_TIME_BASED) {
+
+    executionCtxt->addActiveMetricGroup(metricGroup);
+}
+
+bool SingleMetricStreamerCollector::start() {
+
+    eventPool = ZelloMetricsUtility::createHostVisibleEventPool(executionCtxt->getContextHandle(0), executionCtxt->getDeviceHandle(0));
+    notificationEvent = ZelloMetricsUtility::createHostVisibleEvent(eventPool);
+
+    zet_metric_streamer_desc_t streamerProperties = {};
+    streamerProperties.notifyEveryNReports = notifyReportCount;
+    streamerProperties.samplingPeriod = samplingPeriod;
+    VALIDATECALL(zetMetricStreamerOpen(executionCtxt->getContextHandle(0),
+                                       executionCtxt->getDeviceHandle(0),
+                                       metricGroup,
+                                       &streamerProperties,
+                                       notificationEvent,
+                                       &metricStreamer));
+    //Initial pause
+    std::this_thread::sleep_for(std::chrono::milliseconds(100));
+    return true;
+}
+
+bool SingleMetricStreamerCollector::stop() {
+
+    VALIDATECALL(zetMetricStreamerClose(metricStreamer));
+    VALIDATECALL(zeEventDestroy(notificationEvent));
+    VALIDATECALL(zeEventPoolDestroy(eventPool));
+    return true;
+}
+
+bool SingleMetricStreamerCollector::prefixCommands() {
+    return true;
+}
+
+bool SingleMetricStreamerCollector::suffixCommands() {
+    return true;
+}
+
+bool SingleMetricStreamerCollector::isDataAvailable() {
+    return zeEventQueryStatus(notificationEvent) == ZE_RESULT_SUCCESS;
+}
+
+void SingleMetricStreamerCollector::showResults() {
+
+    // Read raw buffer size.
+    uint32_t maxRawReportCount = maxRequestRawReportCount;
+    size_t rawDataSize = 0;
+    std::vector<uint8_t> rawData{};
+
+    VALIDATECALL(zetMetricStreamerReadData(metricStreamer, maxRawReportCount, &rawDataSize, nullptr));
+    std::cout << "Streamer read requires: " << rawDataSize << " bytes buffer" << std::endl;
+
+    // Read raw data.
+    rawData.resize(rawDataSize, 0);
+    VALIDATECALL(zetMetricStreamerReadData(metricStreamer, maxRawReportCount, &rawDataSize, rawData.data()));
+    std::cout << "Streamer read raw bytes: " << rawDataSize << std::endl;
+
+    ZelloMetricsUtility::obtainCalculatedMetrics(metricGroup, rawData.data(), static_cast<uint32_t>(rawDataSize));
+}
+
+////////////////////////////////
+/////SingleMetricQueryCollector
+////////////////////////////////
+SingleMetricQueryCollector::SingleMetricQueryCollector(ExecutionContext *executionCtxt,
+                                                       const char *metricGroupName) : SingleMetricCollector(executionCtxt, metricGroupName,
+                                                                                                            ZET_METRIC_GROUP_SAMPLING_TYPE_FLAG_EVENT_BASED) {
+
+    executionCtxt->addActiveMetricGroup(metricGroup);
+}
+
+bool SingleMetricQueryCollector::start() {
+
+    eventPool = ZelloMetricsUtility::createHostVisibleEventPool(executionCtxt->getContextHandle(0), executionCtxt->getDeviceHandle(0));
+    notificationEvent = ZelloMetricsUtility::createHostVisibleEvent(eventPool);
+    queryPoolDesc.count = queryPoolCount;
+    queryPoolDesc.type = ZET_METRIC_QUERY_POOL_TYPE_PERFORMANCE;
+
+    // Create metric query pool.
+    VALIDATECALL(zetMetricQueryPoolCreate(executionCtxt->getContextHandle(0), executionCtxt->getDeviceHandle(0), metricGroup, &queryPoolDesc, &queryPoolHandle));
+
+    // Obtain metric query from pool.
+    VALIDATECALL(zetMetricQueryCreate(queryPoolHandle, querySlotIndex, &queryHandle));
+    return true;
+}
+
+bool SingleMetricQueryCollector::stop() {
+
+    VALIDATECALL(zetMetricQueryDestroy(queryHandle));
+    VALIDATECALL(zetMetricQueryPoolDestroy(queryPoolHandle));
+
+    VALIDATECALL(zeEventDestroy(notificationEvent));
+    VALIDATECALL(zeEventPoolDestroy(eventPool));
+    return true;
+}
+
+bool SingleMetricQueryCollector::prefixCommands() {
+    VALIDATECALL(zetCommandListAppendMetricQueryBegin(executionCtxt->getCommandList(0), queryHandle));
+    return true;
+}
+
+bool SingleMetricQueryCollector::suffixCommands() {
+    // Metric query end.
+    VALIDATECALL(zetCommandListAppendMetricQueryEnd(executionCtxt->getCommandList(0), queryHandle, notificationEvent, 0, nullptr));
+
+    // An optional memory barrier to flush gpu caches.
+    VALIDATECALL(zetCommandListAppendMetricMemoryBarrier(executionCtxt->getCommandList(0)));
+    return true;
+}
+
+bool SingleMetricQueryCollector::isDataAvailable() {
+    return zeEventQueryStatus(notificationEvent) == ZE_RESULT_SUCCESS;
+}
+
+void SingleMetricQueryCollector::setAttributes(uint32_t notifyReportCount, zet_metric_query_pool_type_t queryPoolType) {
+    this->notifyReportCount = notifyReportCount;
+    this->queryPoolType = queryPoolType;
+}
+
+void SingleMetricQueryCollector::showResults() {
+
+    size_t rawDataSize = 0;
+    std::vector<uint8_t> rawData{};
+    // Obtain metric query report size.
+    VALIDATECALL(zetMetricQueryGetData(queryHandle, &rawDataSize, nullptr));
+
+    // Obtain report.
+    rawData.resize(rawDataSize);
+    VALIDATECALL(zetMetricQueryGetData(queryHandle, &rawDataSize, rawData.data()));
+
+    ZelloMetricsUtility::obtainCalculatedMetrics(metricGroup, rawData.data(), static_cast<uint32_t>(rawDataSize));
+}

level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics_execution_context.cpp

@@ -0,0 +1,115 @@
+/*
+ * Copyright (C) 2022 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics.h"
+#include "level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics_util.h"
+
+#include <algorithm>
+#include <chrono>
+#include <limits>
+#include <vector>
+
+///////////////////////////
+/////ExecutionContext
+///////////////////////////
+void ExecutionContext::addActiveMetricGroup(zet_metric_group_handle_t &metricGroup) {
+
+    // Add only if not available already
+    if (std::find(activeMetricGroups.begin(), activeMetricGroups.end(), metricGroup) == activeMetricGroups.end()) {
+        activeMetricGroups.push_back(metricGroup);
+    }
+}
+
+void ExecutionContext::removeActiveMetricGroup(zet_metric_group_handle_t metricGroup) {
+
+    if (metricGroup == nullptr) {
+        deactivateMetricGroups();
+        activeMetricGroups.clear();
+        return;
+    }
+
+    // Remove only if available already
+    auto pos = std::find(activeMetricGroups.begin(), activeMetricGroups.end(), metricGroup);
+    if (pos != activeMetricGroups.end()) {
+        activeMetricGroups.erase(pos);
+    }
+}
+
+bool ExecutionContext::activateMetricGroups() {
+    if (activeMetricGroups.size()) {
+        VALIDATECALL(zetContextActivateMetricGroups(getContextHandle(0), getDeviceHandle(0),
+                                                    static_cast<uint32_t>(activeMetricGroups.size()), activeMetricGroups.data()));
+    }
+    return true;
+}
+
+bool ExecutionContext::deactivateMetricGroups() {
+    VALIDATECALL(zetContextActivateMetricGroups(getContextHandle(0), getDeviceHandle(0), 0, nullptr));
+    return true;
+}
+
+//////////////////////////////////////////
+/////SingleDeviceSingleQueueExecutionCtxt
+//////////////////////////////////////////
+bool SingleDeviceSingleQueueExecutionCtxt::initialize(uint32_t deviceIndex, int32_t subDeviceIndex) {
+
+    ZelloMetricsUtility::createL0();
+    driverHandle = ZelloMetricsUtility::getDriver();
+    EXPECT(driverHandle != nullptr);
+    contextHandle = ZelloMetricsUtility::createContext(driverHandle);
+    EXPECT(contextHandle != nullptr);
+
+    deviceHandle = ZelloMetricsUtility::getDevice(driverHandle, deviceIndex);
+    EXPECT(deviceHandle != nullptr);
+    if (subDeviceIndex != -1) {
+        deviceHandle = ZelloMetricsUtility::getSubDevice(deviceHandle, subDeviceIndex);
+        EXPECT(deviceHandle != nullptr);
+    }
+    commandQueue = ZelloMetricsUtility::createCommandQueue(contextHandle, deviceHandle);
+    EXPECT(commandQueue != nullptr);
+    commandList = ZelloMetricsUtility::createCommandList(contextHandle, deviceHandle);
+    EXPECT(commandList != nullptr);
+    return true;
+}
+
+bool SingleDeviceSingleQueueExecutionCtxt::finalize() {
+
+    VALIDATECALL(zeCommandListDestroy(commandList));
+    VALIDATECALL(zeCommandQueueDestroy(commandQueue));
+    VALIDATECALL(zeContextDestroy(contextHandle));
+
+    deviceHandle = nullptr;
+    driverHandle = nullptr;
+    return true;
+}
+
+bool SingleDeviceSingleQueueExecutionCtxt::run() {
+
+    // Close command list.
+    VALIDATECALL(zeCommandListClose(commandList));
+
+    auto startTime = std::chrono::steady_clock::now();
+    uint32_t diff = 0;
+    uint32_t runCount = 0;
+
+    do {
+        // Execute workload.
+        VALIDATECALL(zeCommandQueueExecuteCommandLists(commandQueue, 1, &commandList, nullptr));
+
+        // If using async command queue, explicit sync must be used for correctness.
+        VALIDATECALL(zeCommandQueueSynchronize(commandQueue, std::numeric_limits<uint32_t>::max()));
+
+        auto currentTime = std::chrono::steady_clock::now();
+        diff = static_cast<uint32_t>(std::chrono::duration_cast<std::chrono::milliseconds>(currentTime - startTime).count());
+
+        runCount++;
+    } while (diff < executionTimeInMilliSeconds);
+
+    std::cout << "CommandList Run count : " << runCount << "\n";
+
+    return true;
+}

level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics_util.cpp

@@ -0,0 +1,373 @@
+/*
+ * Copyright (C) 2022 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics_util.h"
+
+#include "level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics.h"
+
+#include <cstring>
+#include <iomanip>
+#include <mutex>
+#include <thread>
+
+namespace ZelloMetricsUtility {
+//////////////////////////
+/// createL0
+///////////////////////////
+void createL0() {
+
+    VALIDATECALL(zeInit(ZE_INIT_FLAG_GPU_ONLY));
+}
+
+ze_driver_handle_t getDriver() {
+    uint32_t driverCount = 0;
+    ze_driver_handle_t driverHandle = {};
+    // Obtain driver.
+    VALIDATECALL(zeDriverGet(&driverCount, nullptr));
+    EXPECT(driverCount > 0);
+    VALIDATECALL(zeDriverGet(&driverCount, &driverHandle));
+    return driverHandle;
+}
+
+ze_context_handle_t createContext(ze_driver_handle_t &driverHandle) {
+    // Obtain context.
+    ze_context_desc_t contextDesc = {ZE_STRUCTURE_TYPE_CONTEXT_DESC, nullptr, 0};
+    ze_context_handle_t contextHandle = {};
+    VALIDATECALL(zeContextCreate(driverHandle, &contextDesc, &contextHandle));
+    return contextHandle;
+}
+
+bool getTestMachineConfiguration(TestMachineConfiguration &machineConfig) {
+
+    std::vector<ze_device_handle_t> devices = {};
+
+    createL0();
+    ze_driver_handle_t driverHandle = getDriver();
+    EXPECT((driverHandle != nullptr));
+
+    uint32_t deviceCount = 0;
+    VALIDATECALL(zeDeviceGet(driverHandle, &deviceCount, nullptr));
+    deviceCount = std::min(deviceCount, MAX_DEVICES_IN_MACHINE);
+    devices.resize(deviceCount);
+    VALIDATECALL(zeDeviceGet(driverHandle, &deviceCount, devices.data()));
+
+    machineConfig.deviceCount = deviceCount;
+    for (uint32_t deviceId = 0; deviceId < deviceCount; deviceId++) {
+        machineConfig.devices[deviceId].deviceIndex = deviceId;
+
+        uint32_t subDevicesCount = 0;
+        VALIDATECALL(zeDeviceGetSubDevices(devices[deviceId], &subDevicesCount, nullptr));
+        machineConfig.devices[deviceId].subDeviceCount = subDevicesCount;
+    }
+    return true;
+}
+
+ze_device_handle_t getDevice(ze_driver_handle_t &driverHandle, uint32_t deviceIndex) {
+
+    uint32_t deviceCount = 0;
+    std::vector<ze_device_handle_t> devices;
+
+    // Obtain all devices.
+    VALIDATECALL(zeDeviceGet(driverHandle, &deviceCount, nullptr));
+    EXPECT((deviceCount > deviceIndex));
+    devices.resize(deviceCount);
+    VALIDATECALL(zeDeviceGet(driverHandle, &deviceCount, devices.data()));
+
+    return devices[deviceIndex];
+}
+
+ze_device_handle_t getSubDevice(ze_device_handle_t &deviceHandle, uint32_t subDeviceIndex) {
+
+    uint32_t subDevicesCount = 0;
+    std::vector<ze_device_handle_t> subDevices = {};
+
+    // Sub devices count.
+    VALIDATECALL(zeDeviceGetSubDevices(deviceHandle, &subDevicesCount, nullptr));
+    EXPECT(subDevicesCount > subDeviceIndex);
+
+    subDevices.resize(subDevicesCount);
+    VALIDATECALL(zeDeviceGetSubDevices(deviceHandle, &subDevicesCount, subDevices.data()));
+    return subDevices[subDeviceIndex];
+}
+
+ze_command_queue_handle_t createCommandQueue(ze_context_handle_t &contextHandle,
+                                             ze_device_handle_t &deviceHandle) {
+
+    uint32_t queueGroupsCount = 0;
+    ze_command_queue_desc_t queueDescription = {};
+
+    VALIDATECALL(zeDeviceGetCommandQueueGroupProperties(deviceHandle, &queueGroupsCount, nullptr));
+
+    if (queueGroupsCount == 0) {
+        std::cout << "No queue groups found!\n";
+        std::terminate();
+    }
+
+    std::vector<ze_command_queue_group_properties_t> queueProperties(queueGroupsCount);
+    VALIDATECALL(zeDeviceGetCommandQueueGroupProperties(deviceHandle, &queueGroupsCount,
+                                                        queueProperties.data()));
+
+    for (uint32_t i = 0; i < queueGroupsCount; ++i) {
+        if (queueProperties[i].flags & ZE_COMMAND_QUEUE_GROUP_PROPERTY_FLAG_COMPUTE) {
+            queueDescription.ordinal = i;
+        }
+    }
+
+    queueDescription.index = 0;
+    queueDescription.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS;
+    ze_command_queue_handle_t commandQueue = {};
+    VALIDATECALL(zeCommandQueueCreate(contextHandle, deviceHandle, &queueDescription, &commandQueue));
+    return commandQueue;
+}
+
+ze_command_list_handle_t createCommandList(ze_context_handle_t &contextHandle, ze_device_handle_t &deviceHandle) {
+
+    ze_command_queue_desc_t queueDescription = {};
+    uint32_t queueGroupsCount = 0;
+    zeDeviceGetCommandQueueGroupProperties(deviceHandle, &queueGroupsCount, nullptr);
+    std::vector<ze_command_queue_group_properties_t> queueProperties(queueGroupsCount);
+    zeDeviceGetCommandQueueGroupProperties(deviceHandle, &queueGroupsCount, queueProperties.data());
+
+    for (uint32_t i = 0; i < queueGroupsCount; ++i) {
+        if (queueProperties[i].flags & ZE_COMMAND_QUEUE_GROUP_PROPERTY_FLAG_COMPUTE) {
+            queueDescription.ordinal = i;
+        }
+    }
+
+    ze_command_list_desc_t commandListDesc = {};
+    commandListDesc.commandQueueGroupOrdinal = queueDescription.ordinal;
+
+    // Create command list.
+    ze_command_list_handle_t commandList = {};
+    VALIDATECALL(zeCommandListCreate(contextHandle, deviceHandle, &commandListDesc, &commandList));
+    return commandList;
+}
+
+void printMetricGroupProperties(const zet_metric_group_properties_t &properties) {
+    std::cout << "METRIC GROUP: "
+              << "name: " << properties.name << ", "
+              << "desc: " << properties.description << ", "
+              << "samplingType: " << properties.samplingType << ", "
+              << "domain: " << properties.domain << ", "
+              << "metricCount: " << properties.metricCount << std::endl;
+}
+
+///////////////////////////
+/// printMetricProperties
+///////////////////////////
+void printMetricProperties(const zet_metric_properties_t &properties) {
+    std::cout << "\tMETRIC: "
+              << "name: " << properties.name << ", "
+              << "desc: " << properties.description << ", "
+              << "component: " << properties.component << ", "
+              << "tier: " << properties.tierNumber << ", "
+              << "metricType: " << properties.metricType << ", "
+              << "resultType: " << properties.resultType << ", "
+              << "units: " << properties.resultUnits << std::endl;
+}
+
+///////////////////////////
+/// wait
+///////////////////////////
+void sleep(uint32_t milliseconds) {
+    std::this_thread::sleep_for(std::chrono::milliseconds(milliseconds));
+}
+
+zet_metric_group_handle_t findMetricGroup(const char *groupName,
+                                          const zet_metric_group_sampling_type_flag_t samplingType,
+                                          ze_device_handle_t deviceHandle) {
+
+    uint32_t metricGroupCount = 0;
+    std::vector<zet_metric_group_handle_t> metricGroups = {};
+
+    // Obtain metric group count for a given device.
+    VALIDATECALL(zetMetricGroupGet(deviceHandle, &metricGroupCount, nullptr));
+
+    // Obtain all metric groups.
+    metricGroups.resize(metricGroupCount);
+    VALIDATECALL(zetMetricGroupGet(deviceHandle, &metricGroupCount, metricGroups.data()));
+
+    // Enumerate metric groups to find a particular one with a given group name
+    // and sampling type requested by the user.
+    for (uint32_t i = 0; i < metricGroupCount; ++i) {
+
+        const zet_metric_group_handle_t metricGroupHandle = metricGroups[i];
+        zet_metric_group_properties_t metricGroupProperties = {};
+
+        // Obtain metric group properties to check the group name and sampling type.
+        VALIDATECALL(zetMetricGroupGetProperties(metricGroupHandle, &metricGroupProperties));
+
+        const bool validGroupName = strcmp(metricGroupProperties.name, groupName) == 0;
+        const bool validSamplingType = (metricGroupProperties.samplingType & samplingType);
+
+        // Validating the name and sampling type.
+        if (validSamplingType) {
+
+            // Obtain metric group handle.
+            if (validGroupName) {
+                return metricGroupHandle;
+            }
+        }
+    }
+
+    // Unable to find metric group.
+    VALIDATECALL(ZE_RESULT_ERROR_UNKNOWN);
+    return nullptr;
+}
+
+ze_event_pool_handle_t createHostVisibleEventPool(ze_context_handle_t contextHandle, ze_device_handle_t deviceHandle) {
+
+    ze_event_pool_handle_t eventPool = {};
+    ze_event_pool_desc_t eventPoolDesc = {};
+    eventPoolDesc.count = 1;
+    eventPoolDesc.flags = ze_event_pool_flag_t::ZE_EVENT_POOL_FLAG_HOST_VISIBLE;
+
+    // Optional notification event to know if Streamer reports are ready to read.
+    VALIDATECALL(zeEventPoolCreate(contextHandle, &eventPoolDesc, 1, &deviceHandle, &eventPool));
+    return eventPool;
+}
+
+ze_event_handle_t createHostVisibleEvent(ze_event_pool_handle_t hostVisibleEventPool) {
+
+    ze_event_desc_t notificationEventDesc = {};
+    ze_event_handle_t notificationEvent = {};
+    notificationEventDesc.index = 0;
+    notificationEventDesc.wait = ze_event_scope_flag_t::ZE_EVENT_SCOPE_FLAG_HOST;
+    notificationEventDesc.signal = ze_event_scope_flag_t::ZE_EVENT_SCOPE_FLAG_DEVICE;
+
+    VALIDATECALL(zeEventCreate(hostVisibleEventPool, &notificationEventDesc, &notificationEvent));
+    return notificationEvent;
+}
+
+void obtainCalculatedMetrics(zet_metric_group_handle_t metricGroup, uint8_t *rawData, uint32_t rawDataSize) {
+
+    uint32_t setCount = 0;
+    uint32_t totalCalculatedMetricCount = 0;
+    zet_metric_group_properties_t properties = {};
+    std::vector<uint32_t> metricCounts = {};
+    std::vector<zet_typed_value_t> results = {};
+    std::vector<zet_metric_handle_t> metrics = {};
+    ze_result_t result = ZE_RESULT_SUCCESS;
+
+    // Obtain maximum space for calculated metrics.
+    result = zetMetricGroupCalculateMetricValues(
+        metricGroup,
+        ZET_METRIC_GROUP_CALCULATION_TYPE_METRIC_VALUES,
+        rawDataSize, rawData,
+        &totalCalculatedMetricCount,
+        nullptr);
+
+    if (result == ZE_RESULT_ERROR_UNKNOWN) {
+
+        // Try to use calculate for multiple metric values.
+        VALIDATECALL(zetMetricGroupCalculateMultipleMetricValuesExp(
+            metricGroup,
+            ZET_METRIC_GROUP_CALCULATION_TYPE_METRIC_VALUES,
+            rawDataSize, rawData,
+            &setCount, &totalCalculatedMetricCount,
+            nullptr, nullptr));
+
+        // Allocate space for calculated reports.
+        metricCounts.resize(setCount);
+        results.resize(totalCalculatedMetricCount);
+
+        // Obtain calculated metrics and their count.
+        VALIDATECALL(zetMetricGroupCalculateMultipleMetricValuesExp(
+            metricGroup,
+            ZET_METRIC_GROUP_CALCULATION_TYPE_METRIC_VALUES,
+            rawDataSize, rawData,
+            &setCount, &totalCalculatedMetricCount,
+            metricCounts.data(), results.data()));
+
+    } else {
+        // Allocate space for calculated reports.
+        setCount = 1;
+        metricCounts.resize(setCount);
+        results.resize(totalCalculatedMetricCount);
+
+        // Obtain calculated metrics and their count.
+        VALIDATECALL(zetMetricGroupCalculateMetricValues(
+            metricGroup,
+            ZET_METRIC_GROUP_CALCULATION_TYPE_METRIC_VALUES,
+            rawDataSize, rawData,
+            &totalCalculatedMetricCount, results.data()));
+
+        metricCounts[0] = totalCalculatedMetricCount;
+    }
+
+    // Obtain metric group properties to show each metric.
+    VALIDATECALL(zetMetricGroupGetProperties(metricGroup, &properties));
+
+    // Allocate space for all metrics from a given metric group.
+    metrics.resize(properties.metricCount);
+
+    // Obtain metrics from a given metric group.
+    VALIDATECALL(zetMetricGet(metricGroup, &properties.metricCount, metrics.data()));
+
+    for (uint32_t i = 0; i < setCount; ++i) {
+
+        std::cout << "\r\nSet " << i;
+
+        const uint32_t metricCount = properties.metricCount;
+        const uint32_t metricCountForSet = metricCounts[i];
+
+        for (uint32_t j = 0; j < metricCountForSet; j++) {
+
+            const uint32_t resultIndex = j + metricCount * i;
+            const uint32_t metricIndex = j % metricCount;
+            zet_metric_properties_t metricProperties = {};
+
+            VALIDATECALL((resultIndex < totalCalculatedMetricCount) ? ZE_RESULT_SUCCESS : ZE_RESULT_ERROR_UNKNOWN)
+
+            // Obtain single metric properties to learn output value type.
+            VALIDATECALL(zetMetricGetProperties(metrics[metricIndex], &metricProperties));
+
+            VALIDATECALL((results[resultIndex].type == metricProperties.resultType) ? ZE_RESULT_SUCCESS : ZE_RESULT_ERROR_UNKNOWN)
+
+            if (metricIndex == 0) {
+                std::cout << "\r\n";
+            }
+
+            std::cout << "\r\n";
+            std::cout << std::setw(25) << metricProperties.name << ": ";
+
+            switch (results[resultIndex].type) {
+            case zet_value_type_t::ZET_VALUE_TYPE_BOOL8:
+                std::cout << std::setw(12);
+                std::cout << (results[resultIndex].value.b8 ? "true" : "false");
+                break;
+
+            case zet_value_type_t::ZET_VALUE_TYPE_FLOAT32:
+                std::cout << std::setw(12);
+                std::cout << results[resultIndex].value.fp32;
+                break;
+
+            case zet_value_type_t::ZET_VALUE_TYPE_FLOAT64:
+                std::cout << std::setw(12);
+                std::cout << results[resultIndex].value.fp64;
+                break;
+
+            case zet_value_type_t::ZET_VALUE_TYPE_UINT32:
+                std::cout << std::setw(12);
+                std::cout << results[resultIndex].value.ui32;
+                break;
+
+            case zet_value_type_t::ZET_VALUE_TYPE_UINT64:
+                std::cout << std::setw(12);
+                std::cout << results[resultIndex].value.ui64;
+                break;
+
+            default:
+                break;
+            }
+        }
+
+        std::cout << "\r\n";
+    }
+}
+
+} // namespace ZelloMetricsUtility

level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics_util.h

@@ -0,0 +1,43 @@
+/*
+ * Copyright (C) 2022 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics.h"
+
+#pragma once
+
+#define MAX_DEVICES_IN_MACHINE (64u)
+
+namespace ZelloMetricsUtility {
+
+// Simplistically assuming one level descendant configuration
+struct TestMachineConfiguration {
+    struct {
+        uint32_t deviceIndex;
+        uint32_t subDeviceCount;
+    } devices[MAX_DEVICES_IN_MACHINE];
+    uint32_t deviceCount = 0;
+};
+
+void createL0();
+ze_driver_handle_t getDriver();
+ze_context_handle_t createContext(ze_driver_handle_t &driverHandle);
+ze_device_handle_t getDevice(ze_driver_handle_t &driverHandle, uint32_t deviceIndex);
+ze_device_handle_t getSubDevice(ze_device_handle_t &deviceHandle, uint32_t subDeviceIndex);
+ze_command_queue_handle_t createCommandQueue(ze_context_handle_t &contextHandle,
+                                             ze_device_handle_t &deviceHandle);
+ze_command_list_handle_t createCommandList(ze_context_handle_t &contextHandle, ze_device_handle_t &deviceHandle);
+void printMetricGroupProperties(const zet_metric_group_properties_t &properties);
+void printMetricProperties(const zet_metric_properties_t &properties);
+void sleep(uint32_t milliseconds);
+bool getTestMachineConfiguration(TestMachineConfiguration &machineConfig);
+zet_metric_group_handle_t findMetricGroup(const char *groupName,
+                                          const zet_metric_group_sampling_type_flag_t samplingType,
+                                          ze_device_handle_t deviceHandle);
+ze_event_pool_handle_t createHostVisibleEventPool(ze_context_handle_t contextHandle, ze_device_handle_t deviceHandle);
+ze_event_handle_t createHostVisibleEvent(ze_event_pool_handle_t hostVisibleEventPool);
+void obtainCalculatedMetrics(zet_metric_group_handle_t metricGroup, uint8_t *rawData, uint32_t rawDataSize);
+} // namespace ZelloMetricsUtility

level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics_workload.cpp

@@ -0,0 +1,197 @@
+/*
+ * Copyright (C) 2022 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "level_zero/tools/test/black_box_tests/zello_metrics/zello_metrics.h"
+
+#include <cstring>
+#include <fstream>
+#include <memory>
+
+///////////////////////////////////////////////////
+/////AppendMemoryCopyFromHeapToDeviceAndBackToHost
+///////////////////////////////////////////////////
+void AppendMemoryCopyFromHeapToDeviceAndBackToHost::initialize() {
+    heapBuffer1 = new char[allocSize];
+    heapBuffer2 = new char[allocSize];
+    for (size_t i = 0; i < allocSize; ++i) {
+        heapBuffer1[i] = static_cast<char>(i + 1);
+    }
+
+    ze_device_mem_alloc_desc_t deviceDesc = {};
+    deviceDesc.stype = ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC;
+    deviceDesc.ordinal = 0;
+    deviceDesc.flags = 0;
+    deviceDesc.pNext = nullptr;
+
+    VALIDATECALL(zeMemAllocDevice(executionCtxt->getContextHandle(0), &deviceDesc, allocSize, allocSize, executionCtxt->getDeviceHandle(0), &zeBuffer));
+}
+
+bool AppendMemoryCopyFromHeapToDeviceAndBackToHost::appendCommands() {
+
+    // Counter to increase complexity of the workload
+    int32_t repeatCount = 3;
+
+    while (repeatCount-- > 0) {
+        // Copy from heap to device-allocated memory
+        VALIDATECALL(zeCommandListAppendMemoryCopy(executionCtxt->getCommandList(0), zeBuffer, heapBuffer1, allocSize,
+                                                   nullptr, 0, nullptr));
+        VALIDATECALL(zeCommandListAppendBarrier(executionCtxt->getCommandList(0), nullptr, 0, nullptr));
+        // Copy from device-allocated memory to heap2
+        VALIDATECALL(zeCommandListAppendMemoryCopy(executionCtxt->getCommandList(0), heapBuffer2, zeBuffer, allocSize,
+                                                   nullptr, 0, nullptr));
+    }
+    return true;
+}
+
+bool AppendMemoryCopyFromHeapToDeviceAndBackToHost::validate() {
+    // Validate heap and ze buffers have the original data from heapBuffer1
+    return (0 == memcmp(heapBuffer1, heapBuffer2, allocSize));
+}
+
+void AppendMemoryCopyFromHeapToDeviceAndBackToHost::finalize() {
+
+    VALIDATECALL(zeMemFree(executionCtxt->getContextHandle(0), zeBuffer));
+    delete[] heapBuffer1;
+    delete[] heapBuffer2;
+}
+
+////////////////////////
+/////CopyBufferToBuffer
+////////////////////////
+void CopyBufferToBuffer::initialize() {
+
+    ze_device_mem_alloc_desc_t deviceDesc = {ZE_STRUCTURE_TYPE_DEVICE_MEM_ALLOC_DESC};
+    deviceDesc.flags = ZE_DEVICE_MEM_ALLOC_FLAG_BIAS_UNCACHED;
+    deviceDesc.ordinal = 0;
+
+    ze_host_mem_alloc_desc_t hostDesc = {ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC};
+    hostDesc.flags = ZE_DEVICE_MEM_ALLOC_FLAG_BIAS_UNCACHED;
+
+    VALIDATECALL(zeMemAllocShared(executionCtxt->getContextHandle(0), &deviceDesc, &hostDesc,
+                                  allocationSize, 1, executionCtxt->getDeviceHandle(0),
+                                  &sourceBuffer));
+    VALIDATECALL(zeMemAllocShared(executionCtxt->getContextHandle(0), &deviceDesc, &hostDesc,
+                                  allocationSize, 1, executionCtxt->getDeviceHandle(0),
+                                  &destinationBuffer));
+
+    // Initialize memory
+    memset(sourceBuffer, 55, allocationSize);
+    memset(destinationBuffer, 0, allocationSize);
+
+    std::ifstream file("copy_buffer_to_buffer.spv", std::ios::binary);
+
+    if (file.is_open()) {
+        file.seekg(0, file.end);
+        auto length = file.tellg();
+        file.seekg(0, file.beg);
+
+        std::cout << "Using copy_buffer_to_buffer.spv" << std::endl;
+
+        std::unique_ptr<char[]> spirvInput(new char[length]);
+        file.read(spirvInput.get(), length);
+
+        ze_module_desc_t moduleDesc = {};
+        ze_module_build_log_handle_t buildlog;
+        moduleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV;
+        moduleDesc.pInputModule = reinterpret_cast<const uint8_t *>(spirvInput.get());
+        moduleDesc.inputSize = length;
+        moduleDesc.pBuildFlags = "";
+
+        if (zeModuleCreate(executionCtxt->getContextHandle(0), executionCtxt->getDeviceHandle(0), &moduleDesc, &module, &buildlog) != ZE_RESULT_SUCCESS) {
+            size_t szLog = 0;
+            zeModuleBuildLogGetString(buildlog, &szLog, nullptr);
+
+            char *strLog = (char *)malloc(szLog);
+            zeModuleBuildLogGetString(buildlog, &szLog, strLog);
+            std::cout << "Build log:" << strLog << std::endl;
+
+            free(strLog);
+        }
+        VALIDATECALL(zeModuleBuildLogDestroy(buildlog));
+
+        ze_kernel_desc_t kernelDesc = {};
+        kernelDesc.pKernelName = "CopyBufferToBufferBytes";
+        VALIDATECALL(zeKernelCreate(module, &kernelDesc, &kernel));
+        file.close();
+        executeFromSpirv = true;
+    }
+}
+
+bool CopyBufferToBuffer::appendCommands() {
+
+    if (executeFromSpirv) {
+        uint32_t groupSizeX = 32u;
+        uint32_t groupSizeY = 1u;
+        uint32_t groupSizeZ = 1u;
+        VALIDATECALL(zeKernelSuggestGroupSize(kernel, allocationSize, 1U, 1U, &groupSizeX, &groupSizeY, &groupSizeZ));
+        VALIDATECALL(zeKernelSetGroupSize(kernel, groupSizeX, groupSizeY, groupSizeZ));
+
+        uint32_t offset = 0;
+        VALIDATECALL(zeKernelSetArgumentValue(kernel, 1, sizeof(destinationBuffer), &destinationBuffer));
+        VALIDATECALL(zeKernelSetArgumentValue(kernel, 0, sizeof(sourceBuffer), &sourceBuffer));
+        VALIDATECALL(zeKernelSetArgumentValue(kernel, 2, sizeof(uint32_t), &offset));
+        VALIDATECALL(zeKernelSetArgumentValue(kernel, 3, sizeof(uint32_t), &offset));
+        VALIDATECALL(zeKernelSetArgumentValue(kernel, 4, sizeof(uint32_t), &offset));
+
+        ze_group_count_t dispatchTraits;
+        dispatchTraits.groupCountX = allocationSize / groupSizeX;
+        dispatchTraits.groupCountY = 1u;
+        dispatchTraits.groupCountZ = 1u;
+
+        VALIDATECALL(zeCommandListAppendLaunchKernel(executionCtxt->getCommandList(0), kernel, &dispatchTraits,
+                                                     nullptr, 0, nullptr));
+    } else {
+
+        std::cout << "Using zeCommandListAppendMemoryCopy" << std::endl;
+
+        VALIDATECALL(zeCommandListAppendMemoryCopy(executionCtxt->getCommandList(0), destinationBuffer,
+                                                   sourceBuffer, allocationSize, nullptr, 0, nullptr));
+    }
+
+    return true;
+}
+
+bool CopyBufferToBuffer::validate() {
+
+    // Validate.
+    const bool outputValidationSuccessful = (memcmp(destinationBuffer, sourceBuffer, allocationSize) == 0);
+
+    if (!outputValidationSuccessful) {
+        // Validate
+        uint8_t *srcCharBuffer = static_cast<uint8_t *>(sourceBuffer);
+        uint8_t *dstCharBuffer = static_cast<uint8_t *>(destinationBuffer);
+        for (size_t i = 0; i < allocationSize; i++) {
+            if (srcCharBuffer[i] != dstCharBuffer[i]) {
+                std::cout << "srcBuffer[" << i << "] = " << static_cast<unsigned int>(srcCharBuffer[i]) << " not equal to "
+                          << "dstBuffer[" << i << "] = " << static_cast<unsigned int>(dstCharBuffer[i]) << "\n";
+                break;
+            }
+        }
+    }
+
+    std::cout << "\nResults validation "
+              << (outputValidationSuccessful ? "PASSED" : "FAILED")
+              << std::endl;
+
+    return outputValidationSuccessful;
+}
+
+void CopyBufferToBuffer::finalize() {
+
+    VALIDATECALL(zeMemFree(executionCtxt->getContextHandle(0), sourceBuffer));
+    VALIDATECALL(zeMemFree(executionCtxt->getContextHandle(0), destinationBuffer));
+
+    if (kernel) {
+        zeKernelDestroy(kernel);
+        kernel = nullptr;
+    }
+
+    if (module) {
+        zeModuleDestroy(module);
+        module = nullptr;
+    }
+}