New interop handler tests for image support in LevelZero

SYCL/Plugin/interop-level-zero-image-get-native-mem.cpp

@@ -0,0 +1,106 @@
+// REQUIRES: level_zero, level_zero_dev_kit
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out
+// RUN:  %GPU_RUN_PLACEHOLDER %t.out 2>&1 | FileCheck %s
+
+// we use the interop to get the native image handle and then use that to make a
+// new image and enumerate the pixels.
+
+// CHECK:      (0 0) -- { 0 0 0 0 }
+// CHECK-NEXT: (1 0) -- { 1 1 1 1 }
+// CHECK-NEXT: (2 0) -- { 2 2 2 2 }
+// CHECK-NEXT: (3 0) -- { 3 3 3 3 }
+// CHECK-NEXT: (0 1) -- { 4 4 4 4 }
+// CHECK-NEXT: (1 1) -- { 5 5 5 5 }
+// CHECK-NEXT: (2 1) -- { 6 6 6 6 }
+// CHECK-NEXT: (3 1) -- { 7 7 7 7 }
+
+// clang++ -fsycl -o las.bin -I$SYCL_HOME/build/install/include/sycl -lze_loader
+// interop-level-zero-image-get-native-mem.cpp
+
+#include <level_zero/ze_api.h>
+#include <sycl.hpp>
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
+using namespace sycl;
+
+int main() {
+#ifdef SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO
+  constexpr auto BE = sycl::backend::ext_oneapi_level_zero;
+  sycl::device D =
+      sycl::ext::oneapi::filter_selector("level_zero:gpu").select_device();
+
+  sycl::context Ctx{D};
+  sycl::queue Q(Ctx, D);
+  auto ZeContext = sycl::get_native<BE>(Ctx);
+  auto ZeDevice = sycl::get_native<BE>(D);
+
+  // -----------  IMAGE STUFF
+  using pixelT = sycl::uint4;        // accessor
+  using ChannelDataT = std::uint8_t; // allocator
+  constexpr long width = 4;
+  constexpr long height = 2;
+  constexpr long numPixels = width * height;
+  ChannelDataT *sourceData =
+      (ChannelDataT *)std::calloc(numPixels * 4, sizeof(ChannelDataT));
+  // initialize data: [ (0 0 0 0)  (1 1 1 1) ...]
+  for (size_t i = 0; i < numPixels; i++) {
+    for (size_t chan = 0; chan < 4; chan++) {
+      size_t idx = (i * 4) + chan;
+      sourceData[idx] = (ChannelDataT)i;
+    }
+  }
+  // 8 bits per channel, four per pixel.
+  sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba;
+  sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8;
+
+  const sycl::range<2> ImgRange_2D(width, height);
+  { // closure
+    // 1 - Create simple image.
+    sycl::image<2> image_2D(sourceData, ChanOrder, ChanType, ImgRange_2D);
+
+    // 2 - Grab it's image handle via the get_native_mem interop.
+    using nativeH = sycl::backend_return_t<BE, sycl::image<2>>;
+    sycl::buffer<nativeH, 1> passBack(range<1>{1});
+
+    Q.submit([&](handler &cgh) {
+       auto image_acc =
+           image_2D.get_access<pixelT, sycl::access::mode::read>(cgh);
+       auto passBackAcc = passBack.get_host_access(sycl::write_only);
+       cgh.host_task([=](const interop_handle &IH) {
+         // There is nothing with image handles in the L0 API except
+         // create and destroy. So let's do that.
+         auto ZeImageH = IH.get_native_mem<BE>(image_acc);
+         passBackAcc[0] = ZeImageH;
+       });
+     }).wait();
+
+    // Now we have the ZeImageH, so let's make a new SYCL image from it.
+    auto passBackAcc = passBack.get_host_access(sycl::read_only);
+    nativeH ZeImageH = passBackAcc[0];
+    sycl::backend_input_t<BE, sycl::image<2>> imageData{
+        ZeImageH, ChanOrder, ChanType, ImgRange_2D,
+        sycl::ext::oneapi::level_zero::ownership::keep};
+    sycl::image<2> NewImg = sycl::make_image<BE, 2>(imageData, Ctx);
+
+    // Then use that image to read and stream out the data.
+    Q.submit([&](handler &cgh) {
+       auto read_acc = NewImg.get_access<pixelT, sycl::access::mode::read>(cgh);
+       sycl::stream out(2024, 400, cgh);
+       cgh.single_task([=]() {
+         for (unsigned y = 0; y < height; y++) {
+           for (unsigned x = 0; x < width; x++) {
+             auto location = sycl::int2{x, y};
+             pixelT somePixel = read_acc.read(location);
+             out << "(" << x << " " << y << ") -- { " << somePixel[0] << " "
+                 << somePixel[1] << " " << somePixel[2] << " " << somePixel[3]
+                 << " }" << sycl::endl;
+           }
+         }
+       });
+     }).wait();
+  } // ~image
+
+#else
+  std::cout << "Missing  Level-Zero backend. Test skipped." << std::endl;
+#endif
+  return 0;
+}

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

@@ -0,0 +1,137 @@
+// REQUIRES: level_zero, level_zero_dev_kit
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out
+// RUN: env ZE_DEBUG=1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 | FileCheck %s
+
+// This test verifies that ownership is working correctly.
+// If ownership is ::transfer then the ~image destructor will end up calling
+// zeImageDestroy
+// CHECK: test  ownership::transfer
+// CHECK: ZE ---> zeImageDestroy
+
+// With ownership ::keep it is must be called manually.
+// CHECK: test  ownership::keep
+// CHECK: zeImageDestroy MANUAL
+
+// No other calls should appear.
+// CHECK-NOT: zeImageDestroy
+
+// clang++ -fsycl -o wfd.bin -I$SYCL_HOME/build/install/include/sycl -lze_loader
+// interop-level-zero-image-ownership.cpp
+
+#include <level_zero/ze_api.h>
+#include <sycl.hpp>
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
+
+using namespace sycl;
+
+void test(sycl::ext::oneapi::level_zero::ownership Ownership) {
+
+  constexpr auto BE = sycl::backend::ext_oneapi_level_zero;
+
+  platform Plt{gpu_selector_v};
+
+  auto Devices = Plt.get_devices();
+
+  if (Devices.size() < 1) {
+    std::cout << "Devices not found" << std::endl;
+    return;
+  }
+
+  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);
+
+  // -----------  Image Fundamentals
+  using pixelT = sycl::uint4;        // accessor
+  using ChannelDataT = std::uint8_t; // allocator
+  sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba;
+  sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8;
+  constexpr uint32_t numChannels = 4; // L0 only supports RGBA at this time.
+
+  constexpr uint32_t width = 8;
+  constexpr uint32_t height = 4;
+  constexpr uint32_t depth = 1;
+
+  const sycl::range<2> ImgRange_2D(width, height);
+
+  // ----------- Basic LevelZero Description
+  ze_image_format_type_t ZeImageFormatType = ZE_IMAGE_FORMAT_TYPE_UINT;
+  size_t ZeImageFormatTypeSize = 8;
+  ze_image_format_layout_t ZeImageFormatLayout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8;
+  ze_image_format_t ZeFormatDesc = {
+      ZeImageFormatLayout,       ZeImageFormatType,
+      ZE_IMAGE_FORMAT_SWIZZLE_R, ZE_IMAGE_FORMAT_SWIZZLE_G,
+      ZE_IMAGE_FORMAT_SWIZZLE_B, ZE_IMAGE_FORMAT_SWIZZLE_A};
+
+  ze_image_desc_t ZeImageDesc_base;
+  ZeImageDesc_base.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC;
+  ZeImageDesc_base.pNext = nullptr;
+  ZeImageDesc_base.flags = ZE_IMAGE_FLAG_KERNEL_WRITE;
+  // ZeImageDesc_base.flags = 0;
+  ZeImageDesc_base.arraylevels = 0;
+  ZeImageDesc_base.miplevels = 0;
+  ZeImageDesc_base.format = ZeFormatDesc;
+
+  {
+    // ------ 2D ------
+    ze_image_desc_t ZeImageDesc_2D = ZeImageDesc_base;
+    ZeImageDesc_2D.type = ZE_IMAGE_TYPE_2D;
+    ZeImageDesc_2D.width = width;
+    ZeImageDesc_2D.height = height;
+    ZeImageDesc_2D.depth = 1;
+
+    ze_image_handle_t ZeHImage_2D;
+    ze_result_t res =
+        zeImageCreate(ZeContext, ZeDevice, &ZeImageDesc_2D, &ZeHImage_2D);
+    if (res != ZE_RESULT_SUCCESS) {
+      std::cout << "unable to create image " << res << std::endl;
+      return;
+    }
+
+    { // closure
+      sycl::backend_input_t<BE, sycl::image<2>> ImageInteropInput_2D{
+          ZeHImage_2D, ChanOrder, ChanType, ImgRange_2D, Ownership};
+      auto Image_2D = sycl::make_image<BE, 2>(ImageInteropInput_2D, Context);
+
+      Queue.submit([&](sycl::handler &cgh) {
+        auto write_acc =
+            Image_2D.get_access<pixelT, sycl::access::mode::write>(cgh);
+
+        cgh.parallel_for(ImgRange_2D, [=](sycl::item<2> Item) {
+          auto location = sycl::int2{Item[0], Item[1]};
+          auto sum = Item[0] + Item[1];
+          const pixelT somepixel = {sum, sum, sum, sum};
+          write_acc.write(location, somepixel);
+        });
+      });
+      Queue.wait_and_throw();
+
+    } // ~image
+    // if ownership was transfer, then the ZeHImage_2D was destroyed as part of
+    // the ~image destruction (or deferred)
+
+    if (Ownership == sycl::ext::oneapi::level_zero::ownership::keep) {
+      zeImageDestroy(ZeHImage_2D);
+      std::cout << "zeImageDestroy MANUAL" << std::endl;
+    }
+
+  } // closure
+}
+
+int main() {
+#ifdef SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO
+  std::cout << "test  ownership::transfer" << std::endl;
+  test(sycl::ext::oneapi::level_zero::ownership::transfer);
+
+  std::cout << "test  ownership::keep" << std::endl;
+  test(sycl::ext::oneapi::level_zero::ownership::keep);
+#else
+  std::cout << "Missing  Level-Zero backend. Test skipped." << std::endl;
+#endif
+  std::cout << "chau" << std::endl;
+  return 0;
+}