[UR][CTS] Add tests for urEnqueueDeviceGlobalRead/Write

test/conformance/device_code/CMakeLists.txt

@@ -24,6 +24,7 @@ macro(add_device_binary SOURCE_FILE)
 endmacro()
 
 add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/bar.cpp)
+add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/device_global.cpp)
 add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fill.cpp)
 add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fill_2d.cpp)
 add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fill_3d.cpp)

test/conformance/device_code/device_global.cpp

@@ -0,0 +1,23 @@
+// Copyright (C) 2023 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>
+
+sycl::ext::oneapi::experimental::device_global<int> dev_var;
+
+int main() {
+
+    sycl::queue deviceQueue;
+    sycl::range<1> numOfItems{1};
+    deviceQueue.submit([&](sycl::handler &cgh) {
+        auto kern = [=](sycl::id<1>) {
+            // just increment
+            dev_var = dev_var + 1;
+        };
+        cgh.parallel_for<class device_global>(numOfItems, kern);
+    });
+
+    return 0;
+}

test/conformance/enqueue/CMakeLists.txt

@@ -4,6 +4,8 @@
 # SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 
 add_conformance_test_with_kernels_environment(enqueue
+    urEnqueueDeviceGlobalVariableRead.cpp
+    urEnqueueDeviceGlobalVariableWrite.cpp
     urEnqueueEventsWait.cpp
     urEnqueueEventsWaitWithBarrier.cpp
     urEnqueueKernelLaunch.cpp

test/conformance/enqueue/urEnqueueDeviceGlobalVariableRead.cpp

@@ -0,0 +1,91 @@
+// Copyright (C) 2023 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 <uur/fixtures.h>
+
+using urEnqueueDeviceGetGlobalVariableReadTest = uur::urGlobalVariableTest;
+UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urEnqueueDeviceGetGlobalVariableReadTest);
+
+TEST_P(urEnqueueDeviceGetGlobalVariableReadTest, Success) {
+
+    ASSERT_SUCCESS(urEnqueueDeviceGlobalVariableWrite(
+        queue, program, global_var.name.c_str(), true, sizeof(global_var.value),
+        0, &global_var.value, 0, nullptr, nullptr));
+
+    size_t global_offset = 0;
+    size_t n_dimensions = 1;
+    size_t global_size = 1;
+
+    // execute the kernel
+    ASSERT_SUCCESS(urEnqueueKernelLaunch(queue, kernel, n_dimensions,
+                                         &global_offset, &global_size, nullptr,
+                                         1, nullptr, nullptr));
+    ASSERT_SUCCESS(urQueueFinish(queue));
+
+    // read global var back to host
+    ASSERT_SUCCESS(urEnqueueDeviceGlobalVariableRead(
+        queue, program, global_var.name.c_str(), true, sizeof(global_var.value),
+        0, &global_var.value, 0, nullptr, nullptr));
+
+    // kernel should increment value
+    ASSERT_EQ(global_var.value, 1);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableReadTest, InvalidNullHandleQueue) {
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableRead(
+                         nullptr, program, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, &global_var.value, 0,
+                         nullptr, nullptr),
+                     UR_RESULT_ERROR_INVALID_NULL_HANDLE);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableReadTest, InvalidNullHandleProgram) {
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableRead(
+                         queue, nullptr, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, &global_var.value, 0,
+                         nullptr, nullptr),
+                     UR_RESULT_ERROR_INVALID_NULL_HANDLE);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableReadTest, InvalidNullPointerName) {
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableRead(
+                         queue, program, nullptr, true,
+                         sizeof(global_var.value), 0, &global_var.value, 0,
+                         nullptr, nullptr),
+                     UR_RESULT_ERROR_INVALID_NULL_POINTER);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableReadTest, InvalidNullPointerDst) {
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableRead(
+                         queue, program, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, nullptr, 0, nullptr,
+                         nullptr),
+                     UR_RESULT_ERROR_INVALID_NULL_POINTER);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableReadTest,
+       InvalidEventWaitListNullEvents) {
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableRead(
+                         queue, program, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, &global_var.value, 1,
+                         nullptr, nullptr),
+                     UR_RESULT_ERROR_INVALID_EVENT_WAIT_LIST);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableReadTest, InvalidEventWaitListZeroSize) {
+    ur_event_handle_t evt = nullptr;
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableRead(
+                         queue, program, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, &global_var.value, 0,
+                         &evt, nullptr),
+                     UR_RESULT_ERROR_INVALID_EVENT_WAIT_LIST);
+}
+TEST_P(urEnqueueDeviceGetGlobalVariableReadTest, InvalidEventWaitInvalidEvent) {
+    ur_event_handle_t inv_evt = nullptr;
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableRead(
+                         queue, program, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, &global_var.value, 1,
+                         &inv_evt, nullptr),
+                     UR_RESULT_ERROR_INVALID_EVENT_WAIT_LIST);
+}

test/conformance/enqueue/urEnqueueDeviceGlobalVariableWrite.cpp

@@ -0,0 +1,68 @@
+// Copyright (C) 2023 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 <uur/fixtures.h>
+
+using urEnqueueDeviceGetGlobalVariableWriteTest = uur::urGlobalVariableTest;
+UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urEnqueueDeviceGetGlobalVariableWriteTest);
+
+TEST_P(urEnqueueDeviceGetGlobalVariableWriteTest, InvalidNullHandleQueue) {
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableWrite(
+                         nullptr, program, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, &global_var.value, 0,
+                         nullptr, nullptr),
+                     UR_RESULT_ERROR_INVALID_NULL_HANDLE);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableWriteTest, InvalidNullHandleProgram) {
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableWrite(
+                         queue, nullptr, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, &global_var.value, 0,
+                         nullptr, nullptr),
+                     UR_RESULT_ERROR_INVALID_NULL_HANDLE);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableWriteTest, InvalidNullPointerName) {
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableWrite(
+                         queue, program, nullptr, true,
+                         sizeof(global_var.value), 0, &global_var.value, 0,
+                         nullptr, nullptr),
+                     UR_RESULT_ERROR_INVALID_NULL_POINTER);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableWriteTest, InvalidNullPointerSrc) {
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableWrite(
+                         queue, program, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, nullptr, 0, nullptr,
+                         nullptr),
+                     UR_RESULT_ERROR_INVALID_NULL_POINTER);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableWriteTest,
+       InvalidEventWaitListNullEvents) {
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableWrite(
+                         queue, program, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, &global_var.value, 1,
+                         nullptr, nullptr),
+                     UR_RESULT_ERROR_INVALID_EVENT_WAIT_LIST);
+}
+
+TEST_P(urEnqueueDeviceGetGlobalVariableWriteTest,
+       InvalidEventWaitListZeroSize) {
+    ur_event_handle_t evt = nullptr;
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableWrite(
+                         queue, program, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, &global_var.value, 0,
+                         &evt, nullptr),
+                     UR_RESULT_ERROR_INVALID_EVENT_WAIT_LIST);
+}
+TEST_P(urEnqueueDeviceGetGlobalVariableWriteTest,
+       InvalidEventWaitInvalidEvent) {
+    ur_event_handle_t inv_evt = nullptr;
+    ASSERT_EQ_RESULT(urEnqueueDeviceGlobalVariableWrite(
+                         queue, program, global_var.name.c_str(), true,
+                         sizeof(global_var.value), 0, &global_var.value, 1,
+                         &inv_evt, nullptr),
+                     UR_RESULT_ERROR_INVALID_EVENT_WAIT_LIST);
+}

test/conformance/testing/include/uur/fixtures.h

@@ -993,6 +993,22 @@ struct urKernelExecutionTest : urKernelTest {
     std::vector<ur_mem_handle_t> buffer_args;
     uint32_t current_arg_index = 0;
 };
+
+template <class T> struct GlobalVar {
+    std::string name;
+    T value;
+};
+
+struct urGlobalVariableTest : uur::urKernelExecutionTest {
+    void SetUp() override {
+        program_name = "device_global";
+        global_var = {"_Z7dev_var", 0};
+        UUR_RETURN_ON_FATAL_FAILURE(uur::urKernelExecutionTest::SetUp());
+    }
+
+    GlobalVar<int> global_var;
+};
+
 } // namespace uur
 
 #endif // UR_CONFORMANCE_INCLUDE_FIXTURES_H_INCLUDED