Merge pull request #1335 from fabiomestre/fabio/queue_e2e_port

[CTS] Add extra tests for Queues

Author: kbenzie

test/conformance/CMakeLists.txt

@@ -135,6 +135,7 @@ if(UR_DPCXX)
     add_subdirectory(kernel)
     add_subdirectory(program)
     add_subdirectory(enqueue)
+    add_subdirectory(integration)
     add_subdirectory(exp_command_buffer)
     add_subdirectory(exp_usm_p2p)
 else()

test/conformance/device_code/CMakeLists.txt

@@ -106,6 +106,9 @@ add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fill_usm.cpp)
 add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/foo.cpp)
 add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/image_copy.cpp)
 add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/mean.cpp)
+add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/cpy_and_mult.cpp)
+add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/cpy_and_mult_usm.cpp)
+add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/multiply.cpp)
 add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/spec_constant.cpp)
 add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/spec_constant_multiple.cpp)
 add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/usm_ll.cpp)

test/conformance/device_code/cpy_and_mult.cpp

@@ -0,0 +1,29 @@
+// Copyright (C) 2024 Intel Corporation
+// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
+// See LICENSE.TXT
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include <CL/sycl.hpp>
+
+int main() {
+    size_t array_size = 16;
+    cl::sycl::queue sycl_queue;
+    std::vector<uint32_t> src(array_size, 1);
+    std::vector<uint32_t> dst(array_size, 1);
+    auto src_buff =
+        cl::sycl::buffer<uint32_t>(src.data(), cl::sycl::range<1>(array_size));
+    auto dst_buff =
+        cl::sycl::buffer<uint32_t>(dst.data(), cl::sycl::range<1>(array_size));
+
+    sycl_queue.submit([&](cl::sycl::handler &cgh) {
+        auto src_acc = src_buff.get_access<cl::sycl::access::mode::read>(cgh);
+        auto dst_acc = dst_buff.get_access<cl::sycl::access::mode::write>(cgh);
+        cgh.parallel_for<class cpy_and_mult>(
+            cl::sycl::range<1>{array_size},
+            [src_acc, dst_acc](cl::sycl::item<1> itemId) {
+                auto id = itemId.get_id(0);
+                dst_acc[id] = src_acc[id] * 2;
+            });
+    });
+    return 0;
+}

test/conformance/device_code/cpy_and_mult_usm.cpp

@@ -0,0 +1,22 @@
+// Copyright (C) 2024 Intel Corporation
+// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
+// See LICENSE.TXT
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include <CL/sycl.hpp>
+
+int main() {
+    size_t array_size = 16;
+    cl::sycl::queue sycl_queue;
+    uint32_t *src = cl::sycl::malloc_device<uint32_t>(array_size, sycl_queue);
+    uint32_t *dst = cl::sycl::malloc_device<uint32_t>(array_size, sycl_queue);
+    sycl_queue.submit([&](cl::sycl::handler &cgh) {
+        cgh.parallel_for<class cpy_and_mult_usm>(
+            cl::sycl::range<1>{array_size},
+            [src, dst](cl::sycl::item<1> itemId) {
+                auto id = itemId.get_id(0);
+                dst[id] = src[id] * 2;
+            });
+    });
+    return 0;
+}

test/conformance/device_code/multiply.cpp

@@ -0,0 +1,20 @@
+// Copyright (C) 2024 Intel Corporation
+// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
+// See LICENSE.TXT
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include <sycl/sycl.hpp>
+
+int main() {
+
+    const size_t inputSize = 1;
+    sycl::queue sycl_queue;
+    uint32_t *inputArray = sycl::malloc_shared<uint32_t>(inputSize, sycl_queue);
+
+    sycl_queue.submit([&](sycl::handler &cgh) {
+        cgh.parallel_for<class MultiplyBy2>(
+            sycl::range<1>(inputSize),
+            [=](sycl::id<1> itemID) { inputArray[itemID] *= 2; });
+    });
+    return 0;
+}

test/conformance/integration/CMakeLists.txt

@@ -0,0 +1,10 @@
+# Copyright (C) 2024 Intel Corporation
+# Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
+# See LICENSE.TXT
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+add_conformance_test_with_kernels_environment(integration
+        QueueEmptyStatus.cpp
+        QueueUSM.cpp
+        QueueBuffer.cpp
+  )

test/conformance/integration/QueueBuffer.cpp

@@ -0,0 +1,108 @@
+// Copyright (C) 2024 Intel Corporation
+// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
+// See LICENSE.TXT
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include "fixtures.h"
+#include <chrono>
+#include <thread>
+
+struct QueueBufferTestWithParam : uur::IntegrationQueueTestWithParam {
+    void SetUp() override {
+        program_name = "cpy_and_mult";
+        UUR_RETURN_ON_FATAL_FAILURE(
+            uur::IntegrationQueueTestWithParam::SetUp());
+    }
+
+    void TearDown() override { uur::IntegrationQueueTestWithParam::TearDown(); }
+
+    void verifyResults(ur_mem_handle_t Buffer, uint32_t ExpectedValue) {
+        uint32_t HostMem[ArraySize] = {};
+        ASSERT_SUCCESS(urEnqueueMemBufferRead(Queue, Buffer, true, 0,
+                                              sizeof(uint32_t) * ArraySize,
+                                              HostMem, 0, nullptr, nullptr));
+
+        for (uint32_t i : HostMem) {
+            ASSERT_EQ(i, ExpectedValue);
+        }
+    }
+
+    ur_mem_handle_t Buffer1 = nullptr;
+    ur_mem_handle_t Buffer2 = nullptr;
+};
+
+UUR_TEST_SUITE_P(QueueBufferTestWithParam,
+                 testing::Values(0, /* In-Order */
+                                 UR_QUEUE_FLAG_OUT_OF_ORDER_EXEC_MODE_ENABLE),
+                 uur::IntegrationQueueTestWithParam::paramPrinter);
+
+/* Submits multiple kernels that interact with each other by accessing and
+ * writing to the same buffers.
+ * Checks that when using an IN_ORDER queue, no synchronization is needed
+ * between calls to urEnqueueKernelLaunch.
+ * Checks that when using an OUT_OF_ORDER queue, synchronizing using only
+ * event barriers is enough. */
+TEST_P(QueueBufferTestWithParam, QueueBufferTest) {
+
+    std::vector<ur_event_handle_t> EventsFill;
+    ur_event_handle_t Event;
+
+    size_t Buffer1Index;
+    size_t Buffer2Index;
+    ASSERT_NO_FATAL_FAILURE(
+        AddBuffer1DArg(ArraySize * sizeof(uint32_t), &Buffer1, &Buffer1Index));
+    ASSERT_NO_FATAL_FAILURE(
+        AddBuffer1DArg(ArraySize * sizeof(uint32_t), &Buffer2, &Buffer2Index));
+
+    ASSERT_SUCCESS(urEnqueueMemBufferFill(
+        Queue, Buffer1, &InitialValue, sizeof(uint32_t), 0,
+        ArraySize * sizeof(uint32_t), 0, nullptr, &Event));
+    EventsFill.push_back(Event);
+
+    ASSERT_SUCCESS(urEnqueueMemBufferFill(
+        Queue, Buffer2, &InitialValue, sizeof(uint32_t), 0,
+        ArraySize * sizeof(uint32_t), 0, nullptr, &Event));
+    EventsFill.push_back(Event);
+
+    ASSERT_NO_FATAL_FAILURE(submitBarrierIfNeeded(EventsFill));
+
+    constexpr size_t GlobalOffset = 0;
+    constexpr size_t NDimensions = 1;
+    constexpr uint32_t NumIterations = 5;
+
+    uint32_t CurValueMem1 = InitialValue;
+    uint32_t CurValueMem2 = InitialValue;
+    for (uint32_t i = 0; i < NumIterations; ++i) {
+
+        /* Copy from DeviceMem1 to DeviceMem2 and multiply by 2 */
+        ASSERT_SUCCESS(
+            urKernelSetArgMemObj(kernel, Buffer2Index, nullptr, Buffer2));
+        ASSERT_SUCCESS(
+            urKernelSetArgMemObj(kernel, Buffer1Index, nullptr, Buffer1));
+
+        ASSERT_SUCCESS(urEnqueueKernelLaunch(Queue, kernel, NDimensions,
+                                             &GlobalOffset, &ArraySize, nullptr,
+                                             0, nullptr, &Event));
+        ASSERT_NO_FATAL_FAILURE(submitBarrierIfNeeded(Event));
+
+        CurValueMem2 = CurValueMem1 * 2;
+
+        /* Copy from DeviceMem1 to DeviceMem2 and multiply by 2 */
+        ASSERT_SUCCESS(
+            urKernelSetArgMemObj(kernel, Buffer1Index, nullptr, Buffer2));
+        ASSERT_SUCCESS(
+            urKernelSetArgMemObj(kernel, Buffer2Index, nullptr, Buffer1));
+
+        ASSERT_SUCCESS(urEnqueueKernelLaunch(Queue, kernel, NDimensions,
+                                             &GlobalOffset, &ArraySize, nullptr,
+                                             0, nullptr, &Event));
+        ASSERT_NO_FATAL_FAILURE(submitBarrierIfNeeded(Event));
+
+        CurValueMem1 = CurValueMem2 * 2;
+    }
+
+    ASSERT_SUCCESS(urQueueFinish(Queue));
+
+    ASSERT_NO_FATAL_FAILURE(verifyResults(Buffer1, CurValueMem1));
+    ASSERT_NO_FATAL_FAILURE(verifyResults(Buffer2, CurValueMem2));
+}

test/conformance/integration/QueueEmptyStatus.cpp

@@ -0,0 +1,107 @@
+// Copyright (C) 2024 Intel Corporation
+// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
+// See LICENSE.TXT
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include "fixtures.h"
+#include <chrono>
+#include <thread>
+
+struct QueueEmptyStatusTestWithParam : uur::IntegrationQueueTestWithParam {
+
+    void SetUp() override {
+
+        program_name = "multiply";
+        UUR_RETURN_ON_FATAL_FAILURE(
+            uur::IntegrationQueueTestWithParam::SetUp());
+
+        ur_device_usm_access_capability_flags_t shared_usm_flags = 0;
+        ASSERT_SUCCESS(
+            uur::GetDeviceUSMSingleSharedSupport(device, shared_usm_flags));
+        if (!(shared_usm_flags & UR_DEVICE_USM_ACCESS_CAPABILITY_FLAG_ACCESS)) {
+            GTEST_SKIP() << "Shared USM is not supported.";
+        }
+
+        ASSERT_SUCCESS(urUSMSharedAlloc(context, device, nullptr, nullptr,
+                                        ArraySize * sizeof(uint32_t),
+                                        &SharedMem));
+    }
+
+    void TearDown() override {
+        ASSERT_SUCCESS(urUSMFree(context, SharedMem));
+        uur::IntegrationQueueTestWithParam::TearDown();
+    }
+
+    void submitWorkToQueue() {
+        ur_event_handle_t Event;
+        ASSERT_SUCCESS(
+            urEnqueueUSMFill(Queue, SharedMem, sizeof(uint32_t), &InitialValue,
+                             ArraySize * sizeof(uint32_t), 0, nullptr, &Event));
+        ASSERT_NO_FATAL_FAILURE(submitBarrierIfNeeded(Event));
+
+        ASSERT_SUCCESS(urKernelSetArgPointer(kernel, 0, nullptr, &SharedMem));
+
+        constexpr size_t global_offset = 0;
+        constexpr size_t n_dimensions = 1;
+        constexpr uint32_t num_iterations = 5;
+        for (uint32_t i = 0; i < num_iterations; ++i) {
+            ASSERT_SUCCESS(urEnqueueKernelLaunch(Queue, kernel, n_dimensions,
+                                                 &global_offset, &ArraySize,
+                                                 nullptr, 0, nullptr, &Event));
+            ASSERT_NO_FATAL_FAILURE(submitBarrierIfNeeded(Event));
+        }
+
+        ASSERT_SUCCESS(urQueueFlush(Queue));
+    }
+
+    void waitUntilQueueEmpty() const {
+
+        using namespace std::chrono_literals;
+
+        constexpr auto step = 500ms;
+        constexpr auto maxWait = 5000ms;
+
+        /* Wait a bit until work finishes running. We don't synchronize with
+         * urQueueFinish() because we want to check if the status is set without
+         * calling it explicitly. */
+        for (auto currentWait = 0ms; currentWait < maxWait;
+             currentWait += step) {
+            std::this_thread::sleep_for(step);
+
+            ur_bool_t is_queue_empty;
+            ASSERT_SUCCESS(urQueueGetInfo(Queue, UR_QUEUE_INFO_EMPTY,
+                                          sizeof(ur_bool_t), &is_queue_empty,
+                                          nullptr));
+            if (is_queue_empty) {
+                return;
+            }
+        }
+
+        /* If we are here, the test failed. Let's call queue finish to avoid
+         * issues when freeing memory */
+        ASSERT_SUCCESS(urQueueFinish(Queue));
+        GTEST_FAIL();
+    }
+
+    void *SharedMem = nullptr;
+};
+
+UUR_TEST_SUITE_P(QueueEmptyStatusTestWithParam,
+                 testing::Values(0, /* In-Order */
+                                 UR_QUEUE_FLAG_OUT_OF_ORDER_EXEC_MODE_ENABLE),
+                 uur::IntegrationQueueTestWithParam::paramPrinter);
+
+/* Submits kernels that have a dependency on each other and checks that the
+ * queue submits all the work in the correct order to the device.
+ * Explicit synchronization (except for barriers) is avoided in these tests to
+ * check that the properties of In-Order and OutOfOrder queues are working as
+ * expected */
+TEST_P(QueueEmptyStatusTestWithParam, QueueEmptyStatusTest) {
+    ASSERT_NO_FATAL_FAILURE(submitWorkToQueue());
+    ASSERT_NO_FATAL_FAILURE(waitUntilQueueEmpty());
+
+    constexpr size_t expected_value = 3200;
+    for (uint32_t i = 0; i < ArraySize; ++i) {
+        ASSERT_EQ(reinterpret_cast<uint32_t *>(SharedMem)[i], expected_value);
+    }
+}