[SYCL][Bindless][4/4] Add tests for experimental implementation of SYCL bindless images extension

sycl/test-e2e/CMakeLists.txt

@@ -24,6 +24,8 @@ endif() # Standalone.
 find_package(Threads REQUIRED)
 set(SYCL_THREADS_LIB ${CMAKE_THREAD_LIBS_INIT})
 
+find_package(Vulkan)
+
 if(NOT LLVM_LIT)
   find_program(LLVM_LIT
                NAMES llvm-lit lit.py lit

sycl/test-e2e/bindless_images/image_get_info.cpp

@@ -0,0 +1,204 @@
+// REQUIRES: linux
+// REQUIRES: cuda
+
+// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
+// RUN: %t.out
+
+#include <iostream>
+#include <sycl/sycl.hpp>
+
+// Uncomment to print additional test information
+// #define VERBOSE_PRINT
+
+void printString(std::string name) {
+#ifdef VERBOSE_PRINT
+  std::cout << name;
+#endif
+}
+
+int main() {
+
+  sycl::device dev;
+  sycl::queue q(dev);
+  auto ctxt = q.get_context();
+
+  size_t height = 13;
+  size_t width = 7;
+  size_t depth = 11;
+
+  bool validated = true;
+
+  try {
+    // Submit dummy kernel to let the runtime decide the backend (CUDA)
+    // Without this, the default Level Zero backend is active
+    q.submit([&](sycl::handler &cgh) { cgh.single_task([]() {}); });
+
+    // Extension: image descriptor - can use the same for both images
+    sycl::ext::oneapi::experimental::image_descriptor desc(
+        {width, height, depth}, sycl::image_channel_order::r,
+        sycl::image_channel_type::signed_int32);
+
+    // Extension: returns the device pointer to the allocated memory
+    // Input images memory
+    sycl::ext::oneapi::experimental::image_mem imgMem(desc, dev, ctxt);
+
+    // Extension: query for bindless image support -- device aspects
+    bool bindlessSupport = dev.has(sycl::aspect::ext_oneapi_bindless_images);
+    bool bindlessSharedUsmSupport =
+        dev.has(sycl::aspect::ext_oneapi_bindless_images_shared_usm);
+    bool usm1dSupport =
+        dev.has(sycl::aspect::ext_oneapi_bindless_images_1d_usm);
+    bool usm2dSupport =
+        dev.has(sycl::aspect::ext_oneapi_bindless_images_2d_usm);
+
+#ifdef VERBOSE_PRINT
+    std::cout << "bindless_images_support: " << bindlessSupport
+              << "\nbindless_images_shared_usm_support: "
+              << bindlessSharedUsmSupport
+              << "\nbindless_images_1d_usm_support: " 1dS
+              << "\nbindless_images_2d_usm_support: " << S << "\n";
+#endif
+
+    // Extension: get pitch alignment information from device -- device info
+    // Make sure our pitch alignment queries work properly
+    // These can be different depending on the device so we cannot test that the
+    // values are correct
+    // But we should at least see that the query itself works
+    auto pitchAlign = dev.get_info<
+        sycl::ext::oneapi::experimental::info::device::image_row_pitch_align>();
+    auto maxPitch = dev.get_info<sycl::ext::oneapi::experimental::info::device::
+                                     max_image_linear_row_pitch>();
+    auto maxWidth = dev.get_info<sycl::ext::oneapi::experimental::info::device::
+                                     max_image_linear_width>();
+    auto maxheight = dev.get_info<sycl::ext::oneapi::experimental::info::
+                                      device::max_image_linear_height>();
+
+#ifdef VERBOSE_PRINT
+    std::cout << "image_row_pitch_align: " << pitchAlign
+              << "\nmax_image_linear_row_pitch: " << maxPitch
+              << "\nmax_image_linear_width: " << maxWidth
+              << "\nmax_image_linear_height: " << maxheight << "\n";
+#endif
+
+    // Extension: query for bindless image mipmaps support -- aspects & info
+    bool mipmapSupport = dev.has(sycl::aspect::ext_oneapi_mipmap);
+    bool mipmapAnisotropySupport =
+        dev.has(sycl::aspect::ext_oneapi_mipmap_anisotropy);
+    float mipmapMaxAnisotropy = dev.get_info<
+        sycl::ext::oneapi::experimental::info::device::mipmap_max_anisotropy>();
+    bool mipmapLevelReferenceSupport =
+        dev.has(sycl::aspect::ext_oneapi_mipmap_level_reference);
+
+#ifdef VERBOSE_PRINT
+    std::cout << "mipmapSupport: " << mipmapSupport
+              << "\nmipmapAnisotropySupport: " << mipmapAnisotropySupport
+              << "\nmipmapMaxAnisotropy: " << mipmapMaxAnisotropy
+              << "\nmipmapLevelReferenceSupport: "
+              << mipmapLevelReferenceSupport << "\n";
+#endif
+
+    // Extension: query for bindless image interop support -- device aspects
+    bool interopMemoryImportSupport =
+        dev.has(sycl::aspect::ext_oneapi_interop_memory_import);
+    bool interopMemoryExportSupport =
+        dev.has(sycl::aspect::ext_oneapi_interop_memory_export);
+    bool interopSemaphoreImportSupport =
+        dev.has(sycl::aspect::ext_oneapi_interop_semaphore_import);
+    bool interopSemaphoreExportSupport =
+        dev.has(sycl::aspect::ext_oneapi_interop_semaphore_export);
+
+#ifdef VERBOSE_PRINT
+    std::cout << "interopMemoryImportSupport: " << interopMemoryImportSupport
+              << "\ninteropMemoryExportSupport: " << interopMemoryExportSupport
+              << "\ninteropSemaphoreImportSupport: "
+              << interopSemaphoreImportSupport
+              << "\ninteropSemaphoreExportSupport: "
+              << interopSemaphoreExportSupport << "\n";
+#endif
+
+    auto rangeMem = imgMem.get_range();
+    auto range = sycl::ext::oneapi::experimental::get_image_range(
+        imgMem.get_handle(), dev, ctxt);
+    if (rangeMem != range) {
+      printString("handle and mem object disagree on image dimensions!\n");
+      validated = false;
+    }
+    if (range[0] == width) {
+      printString("width is correct!\n");
+    } else {
+      printString("width is NOT correct!\n");
+      validated = false;
+    }
+    if (range[1] == height) {
+      printString("height is correct!\n");
+    } else {
+      printString("height is NOT correct!\n");
+      validated = false;
+    }
+    if (range[2] == depth) {
+      printString("depth is correct!\n");
+    } else {
+      printString("depth is NOT correct!\n");
+      validated = false;
+    }
+
+    auto type = imgMem.get_type();
+    if (type == sycl::ext::oneapi::experimental::image_type::standard) {
+      printString("image type is correct!\n");
+    } else {
+      printString("image type is NOT correct!\n");
+      validated = false;
+    }
+
+    auto ctypeMem = imgMem.get_channel_type();
+    auto ctype = sycl::ext::oneapi::experimental::get_image_channel_type(
+        imgMem.get_handle(), dev, ctxt);
+    if (ctypeMem != ctype) {
+      printString("handle and mem object disagree on image channel type!\n");
+      validated = false;
+    }
+    if (ctype == sycl::image_channel_type::signed_int32) {
+      printString("channel type is correct!\n");
+    } else {
+      printString("channel type is NOT correct!\n");
+      validated = false;
+    }
+
+    auto corder = imgMem.get_channel_order();
+    if (corder == sycl::image_channel_order::r) {
+      printString("channel order is correct!\n");
+    } else {
+      printString("channel order is NOT correct!\n");
+      validated = false;
+    }
+
+    auto numchannelsMem = imgMem.get_num_channels();
+    auto numchannels = sycl::ext::oneapi::experimental::get_image_num_channels(
+        imgMem.get_handle(), dev, ctxt);
+    if (numchannelsMem != numchannels) {
+      printString("handle and mem object disagree on number of channels!\n");
+      validated = false;
+    }
+    if (numchannels == 1) {
+      printString("num channels is correct!\n");
+    } else {
+      printString("num channels is NOT correct!\n");
+      validated = false;
+    }
+
+  } catch (sycl::exception e) {
+    std::cerr << "SYCL exception caught! : " << e.what() << "\n";
+    return 1;
+  } catch (...) {
+    std::cerr << "Unknown exception caught!\n";
+    return 2;
+  }
+
+  if (validated) {
+    std::cout << "Test Passed!\n";
+    return 0;
+  }
+
+  std::cout << "Test Failed!" << std::endl;
+  return 3;
+}

sycl/test-e2e/bindless_images/mipmap/mipmap_read_1D.cpp

@@ -0,0 +1,142 @@
+// REQUIRES: linux
+// REQUIRES: cuda
+
+// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
+// RUN: %t.out
+
+#include <iostream>
+#include <sycl/sycl.hpp>
+
+// Uncomment to print additional test information
+// #define VERBOSE_PRINT
+
+class image_addition;
+
+int main() {
+
+  sycl::device dev;
+  sycl::queue q(dev);
+  auto ctxt = q.get_context();
+
+  // declare image data
+  constexpr size_t N = 16;
+  std::vector<float> out(N);
+  std::vector<float> expected(N);
+  std::vector<sycl::float4> dataIn1(N);
+  std::vector<sycl::float4> dataIn2(N / 2);
+  std::vector<sycl::float4> copyOut(N / 2);
+  int j = 0;
+  for (int i = 0; i < N; i++) {
+    expected[i] = i + (j + 10);
+    if (i % 2)
+      j++;
+    dataIn1[i] = sycl::float4(i, i, i, i);
+    if (i < (N / 2)) {
+      dataIn2[i] = sycl::float4(i + 10, i + 10, i + 10, i + 10);
+      copyOut[i] = sycl::float4{0, 0, 0, 0};
+    }
+  }
+
+  try {
+
+    size_t width = N;
+    unsigned int numLevels = 2;
+
+    // Extension: image descriptor -- number of levels
+    sycl::ext::oneapi::experimental::image_descriptor desc(
+        {width}, sycl::image_channel_order::rgba,
+        sycl::image_channel_type::fp32,
+        sycl::ext::oneapi::experimental::image_type::mipmap, numLevels);
+
+    // Extension: allocate mipmap memory on device
+    sycl::ext::oneapi::experimental::image_mem mipMem(desc, dev, ctxt);
+
+    // Extension: retrieve level 0
+    sycl::ext::oneapi::experimental::image_mem_handle imgMem1 =
+        mipMem.get_mip_level_mem_handle(0);
+
+    // Extension: copy over data to device at level 0
+    q.ext_oneapi_copy(dataIn1.data(), imgMem1, desc);
+
+    // Extension: copy data to device at level 1
+    q.ext_oneapi_copy(dataIn2.data(), mipMem.get_mip_level_mem_handle(1),
+                      desc.get_mip_level_desc(1));
+    q.wait_and_throw();
+
+    // Extension: define a sampler object -- extended mipmap attributes
+    sycl::ext::oneapi::experimental::bindless_image_sampler samp(
+        sycl::addressing_mode::mirrored_repeat,
+        sycl::coordinate_normalization_mode::normalized,
+        sycl::filtering_mode::nearest, sycl::filtering_mode::nearest, 0.0f,
+        (float)numLevels, 8.0f);
+
+    // Extension: create a sampled image handle to represent the mipmap
+    sycl::ext::oneapi::experimental::sampled_image_handle mipHandle =
+        sycl::ext::oneapi::experimental::create_image(mipMem, samp, desc, dev,
+                                                      ctxt);
+
+    sycl::buffer<float, 1> buf((float *)out.data(), N);
+    q.submit([&](sycl::handler &cgh) {
+      auto outAcc = buf.get_access<sycl::access_mode::write>(cgh, N);
+
+      cgh.parallel_for<image_addition>(N, [=](sycl::id<1> id) {
+        float sum = 0;
+        float x = float(id[0] + 0.5) / (float)N;
+        // Extension: read mipmap level 0 with anisotropic filtering and level 1
+        // with LOD
+        sycl::float4 px1 =
+            sycl::ext::oneapi::experimental::read_image<sycl::float4>(
+                mipHandle, x, 0.0f, 0.0f);
+        sycl::float4 px2 =
+            sycl::ext::oneapi::experimental::read_image<sycl::float4>(mipHandle,
+                                                                      x, 1.0f);
+
+        sum = px1[0] + px2[0];
+        outAcc[id] = sum;
+      });
+    });
+
+    q.wait_and_throw();
+
+    // Extension: copy data from device
+    q.ext_oneapi_copy(mipMem.get_mip_level_mem_handle(1), copyOut.data(),
+                      desc.get_mip_level_desc(1));
+    q.wait_and_throw();
+
+    // Extension: cleanup
+    sycl::ext::oneapi::experimental::destroy_image_handle(mipHandle, dev, ctxt);
+
+  } catch (sycl::exception e) {
+    std::cerr << "SYCL exception caught! : " << e.what() << "\n";
+    return 1;
+  } catch (...) {
+    std::cerr << "Unknown exception caught!\n";
+    return 2;
+  }
+
+  // collect and validate output
+  bool validated = true;
+  for (int i = 0; i < N; i++) {
+    bool mismatch = false;
+    if (out[i] != expected[i]) {
+      mismatch = true;
+      validated = false;
+    }
+
+    if (mismatch) {
+#ifdef VERBOSE_PRINT
+      std::cout << "Result mismatch! Expected: " << expected[i]
+                << ", Actual: " << out[i] << std::endl;
+#else
+      break;
+#endif
+    }
+  }
+  if (validated) {
+    std::cout << "Test passed!" << std::endl;
+    return 0;
+  }
+
+  std::cout << "Test failed!" << std::endl;
+  return 3;
+}