[SYCL] Test buffer interop for the Level Zero backend

SYCL/Plugin/interop-level-zero-buffer-helpers.hpp

@@ -0,0 +1,18 @@
+#pragma once
+#include "interop-level-zero-buffer-helpers.hpp"
+#include <CL/sycl.hpp>
+// clang-format off
+#include <level_zero/ze_api.h>
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
+// clang-format on
+
+using namespace sycl;
+
+bool is_discrete(const device &Device) {
+  auto ZeDevice = get_native<backend::ext_oneapi_level_zero>(Device);
+  ze_device_properties_t ZeDeviceProps;
+  ZeDeviceProps.stype = ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES;
+  ZeDeviceProps.pNext = nullptr;
+  zeDeviceGetProperties(ZeDevice, &ZeDeviceProps);
+  return !(ZeDeviceProps.flags & ZE_DEVICE_PROPERTY_FLAG_INTEGRATED);
+}

SYCL/Plugin/interop-level-zero-buffer-multi-dim.cpp

@@ -0,0 +1,123 @@
+// REQUIRES: level_zero, level_zero_dev_kit
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out
+// RUN: env SYCL_DEVICE_FILTER=level_zero %GPU_RUN_PLACEHOLDER %t.out
+
+// Test 2D and 3D interoperability buffers for the Level Zero backend.
+
+#include "interop-level-zero-buffer-helpers.hpp"
+#include <CL/sycl.hpp>
+// clang-format off
+#include <level_zero/ze_api.h>
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
+// clang-format on
+
+using namespace cl::sycl;
+
+int main() {
+#ifdef SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO
+  try {
+    platform Plt{gpu_selector{}};
+
+    auto Devices = Plt.get_devices();
+
+    if (Devices.size() < 1) {
+      std::cout << "Devices not found" << std::endl;
+      return 0;
+    }
+
+    device Device = Devices[0];
+    context Context{Device};
+    queue Queue{Context, Device};
+
+    // Get native Level Zero handles
+    auto ZeContext = get_native<backend::ext_oneapi_level_zero>(Context);
+    auto ZeDevice = get_native<backend::ext_oneapi_level_zero>(Device);
+
+    ze_host_mem_alloc_desc_t HostDesc = {};
+    HostDesc.stype = ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC;
+    HostDesc.pNext = nullptr;
+    HostDesc.flags = 0;
+
+    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;
+
+    // Test case #1
+    // Check 2D buffer
+    void *NativeBuffer = nullptr;
+    if (is_discrete(Device))
+      zeMemAllocDevice(ZeContext, &DeviceDesc, 24 * sizeof(int), 1, ZeDevice,
+                       &NativeBuffer);
+    else
+      zeMemAllocHost(ZeContext, &HostDesc, 24 * sizeof(int), 1, &NativeBuffer);
+    {
+      backend_input_t<backend::ext_oneapi_level_zero, buffer<int, 1>>
+          BufferInteropInput = {NativeBuffer,
+                                ext::oneapi::level_zero::ownership::keep};
+      auto BufferInterop = make_buffer<backend::ext_oneapi_level_zero, int, 1>(
+          BufferInteropInput, Context);
+
+      auto Buf2D = BufferInterop.reinterpret<int>(range<2>(4, 6));
+
+      Queue.submit([&](cl::sycl::handler &CGH) {
+        auto Acc2D = Buf2D.get_access<cl::sycl::access::mode::read_write>(CGH);
+        CGH.single_task<class SimpleKernel2D>([=]() {
+          for (int i = 0; i < 4; i++) {
+            for (int j = 0; j < 6; j++) {
+              Acc2D[i][j] = i + j;
+            }
+          }
+        });
+      });
+      Queue.wait();
+
+      {
+        auto HostAcc2D = Buf2D.get_host_access();
+        for (int i = 0; i < 4; i++) {
+          for (int j = 0; j < 6; j++) {
+            assert(HostAcc2D[i][j] == i + j);
+          }
+        }
+      }
+
+      // Test case #2
+      // Check 3D buffer
+      auto Buf3D = BufferInterop.reinterpret<int>(range<3>(4, 2, 3));
+
+      Queue.submit([&](cl::sycl::handler &CGH) {
+        auto Acc3D = Buf3D.get_access<cl::sycl::access::mode::read_write>(CGH);
+        CGH.single_task<class SimpleKernel3D>([=]() {
+          for (int i = 0; i < 4; i++) {
+            for (int j = 0; j < 2; j++) {
+              for (int k = 0; k < 3; k++) {
+                Acc3D[i][j][k] = i + j + k;
+              }
+            }
+          }
+        });
+      });
+      Queue.wait();
+      {
+        auto HostAcc3D = Buf3D.get_host_access();
+        for (int i = 0; i < 4; i++) {
+          for (int j = 0; j < 2; j++) {
+            for (int k = 0; k < 3; k++) {
+              assert(HostAcc3D[i][j][k] == i + j + k);
+            }
+          }
+        }
+      }
+    }
+    zeMemFree(ZeContext, NativeBuffer);
+  } catch (sycl::exception &e) {
+    std::cout << e.what() << std::endl;
+    return 0;
+  }
+#else
+  std::cout << "Test skipped due to missing support for Level-Zero backend."
+            << std::endl;
+#endif
+  return 0;
+}

SYCL/Plugin/interop-level-zero-buffer-ownership.cpp

@@ -0,0 +1,138 @@
+// REQUIRES: level_zero, level_zero_dev_kit
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out
+// RUN: env SYCL_DEVICE_FILTER=level_zero ZE_DEBUG=1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 | FileCheck %s
+
+// Test for Level Zero buffer interop API.
+// Check the following cases:
+// 1. User-provided memory allocation is not freed by DPCPP RT if
+// "keep" ownership is specified.
+// 2. User-provided memory allocation is freed by DPCPP RT if
+// "transfer" ownership is specified.
+
+// Keep ownership
+// CHECK: zeMemFree
+
+// Transfer ownership
+// CHECK: zeMemFree
+// CHECK: zeMemFree
+
+// No other calls to zeMemFree
+// CHECK-NOT: zeMemFree
+
+#include "interop-level-zero-buffer-helpers.hpp"
+#include <CL/sycl.hpp>
+// clang-format off
+#include <level_zero/ze_api.h>
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
+// clang-format on
+
+using namespace cl::sycl;
+
+// Test copy back depending on provided ownership and check that memory is freed
+// properly.
+void test_copyback_and_free(
+    queue &Queue1, queue &Queue2,
+    const ext::oneapi::level_zero::ownership &Ownership) {
+  try {
+    auto Context = Queue1.get_context();
+    auto Device = Queue1.get_info<info::queue::device>();
+
+    // Get native Level Zero handles
+    auto ZeContext = get_native<backend::ext_oneapi_level_zero>(Context);
+    auto ZeDevice = get_native<backend::ext_oneapi_level_zero>(Device);
+
+    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;
+
+    ze_host_mem_alloc_desc_t HostDesc = {};
+    HostDesc.stype = ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC;
+    HostDesc.pNext = nullptr;
+    HostDesc.flags = 0;
+
+    void *NativeBuffer = nullptr;
+    if (is_discrete(Device))
+      // Use shared allocation (the check contents on the host later) associated
+      // with a device if the device is discreet.
+      zeMemAllocShared(ZeContext, &DeviceDesc, &HostDesc, 12 * sizeof(int), 1,
+                       ZeDevice, &NativeBuffer);
+    else
+      // Use host allocation if device is integrated.
+      zeMemAllocHost(ZeContext, &HostDesc, 12 * sizeof(int), 1, &NativeBuffer);
+
+    backend_input_t<backend::ext_oneapi_level_zero, buffer<int, 1>>
+        BufferInteropInput = {NativeBuffer, Ownership};
+    {
+      auto BufferInterop = make_buffer<backend::ext_oneapi_level_zero, int, 1>(
+          BufferInteropInput, Context);
+
+      auto Event = Queue1.submit([&](cl::sycl::handler &CGH) {
+        auto Acc =
+            BufferInterop.get_access<cl::sycl::access::mode::read_write>(CGH);
+        CGH.single_task<class SimpleKernel6>([=]() {
+          for (int i = 0; i < 12; i++) {
+            Acc[i] = 99;
+          }
+        });
+      });
+      Event.wait();
+
+      // Submit in a different context
+      Queue2.submit([&](cl::sycl::handler &CGH) {
+        auto Acc =
+            BufferInterop.get_access<cl::sycl::access::mode::read_write>(CGH);
+        CGH.single_task<class SimpleKernel7>([=]() {
+          for (int i = 0; i < 12; i++) {
+            Acc[i] *= 2;
+          }
+        });
+      });
+
+      Queue2.wait();
+    }
+    if (Ownership == ext::oneapi::level_zero::ownership::keep)
+      zeMemFree(ZeContext, NativeBuffer);
+  } catch (sycl::exception &e) {
+    std::cout << e.what() << std::endl;
+  }
+}
+
+int main() {
+#ifdef SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO
+  try {
+    platform Plt{gpu_selector{}};
+
+    auto Devices = Plt.get_devices();
+
+    if (Devices.size() < 1) {
+      std::cout << "Devices not found" << std::endl;
+      return 0;
+    }
+
+    device Dev1 = Devices[0];
+    context Context1{Dev1};
+    queue Queue1{Context1, Dev1};
+
+    device Dev2 = Devices.size() > 1 ? Devices[1] : Devices[0];
+    context Context2{Dev2};
+    queue Queue2{Context2, Dev2};
+
+    std::cout << "Test case #1: Keep ownership" << std::endl;
+    test_copyback_and_free(Queue1, Queue2,
+                           ext::oneapi::level_zero::ownership::keep);
+
+    std::cout << "Test case #2: Transfer ownership" << std::endl;
+    test_copyback_and_free(Queue1, Queue2,
+                           ext::oneapi::level_zero::ownership::transfer);
+
+  } catch (exception &e) {
+    std::cout << e.what() << std::endl;
+  }
+#else
+  std::cout << "Test skipped due to missing support for Level-Zero backend."
+            << std::endl;
+#endif
+  return 0;
+}