Merge remote-tracking branch 'origin/sycl' into no-unsupported

Author: KornevNikita

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -241,7 +241,9 @@ descriptor against the limitations outlined below. If the given descriptor is
 deemed invalid, then a `sycl::exception` will be thrown with error code 
 `sycl::errc::invalid`.
 
-For all image types, the value of `num_channels` must be `1`, `2`, or `4`.
+The value of `num_channels` supported by all image types and backends
+is  `1`, `2`, or `4`.
+Some backends also support `num_channels` to be `3`.
 
 For the `standard` image type, the value of `num_levels` and `array_size` must
 both be `1`.
@@ -2107,6 +2109,12 @@ There are dimension specific limitations:
 * 3D - No support at the moment. Possible support in non CUDA backends in the 
        future.
 
+=== 3 channel format support
+
+The ability to create an image with 3 channels depends on the backend.
+There is currently no way to query a backend whether it supports this feature.
+This query should be added in a later revision of the proposal.
+
 === Not supported yet
 
 These features still need to be handled:
@@ -2319,4 +2327,5 @@ These features still need to be handled:
                    sub-region copies.
 |6.4|2024-10-15| - Fix bindless spec examples and include examples in bindless
                    spec using asciidoc include.
+|6.5|2024-10-22| - Allow 3-channel image formats on some backends.
 |======================

sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp

@@ -23,13 +23,15 @@ namespace ext::oneapi::experimental {
 
 namespace detail {
 
-inline image_channel_order
+constexpr image_channel_order
 get_image_default_channel_order(unsigned int num_channels) {
   switch (num_channels) {
   case 1:
     return image_channel_order::r;
   case 2:
     return image_channel_order::rg;
+  case 3:
+    return image_channel_order::rgb;
   case 4:
     return image_channel_order::rgba;
   default:
@@ -120,13 +122,9 @@ struct image_descriptor {
   }
 
   void verify() const {
-
-    if (this->num_channels != 1 && this->num_channels != 2 &&
-        this->num_channels != 4) {
-      // Images can only have 1, 2, or 4 channels.
+    if ((this->num_channels < 1) || (this->num_channels > 4)) {
       throw sycl::exception(sycl::errc::invalid,
-                            "Images must have only 1, 2, or 4 channels! Use a "
-                            "valid number of channels instead.");
+                            "Images must have 1, 2, 3, or 4 channels.");
     }
 
     switch (this->type) {

sycl/test-e2e/Basic/interop/get_native_ocl.cpp

@@ -1,5 +1,5 @@
-// REQUIRES: opencl, opencl_dev_kit
-// RUN: %{build} %opencl_options -o %t.ocl.out
+// REQUIRES: opencl, opencl_icd
+// RUN: %{build} %opencl_lib -o %t.out
 // RUN: %{run} %t.out
 
 #include <CL/cl.h>

sycl/test-e2e/Basic/interop/interop_all_backends.cpp

@@ -1,25 +1,29 @@
-// REQUIRES: CUDA || HIP
-// RUN: %{build} %if hip %{ -DSYCL_EXT_ONEAPI_BACKEND_HIP %} %else %{ %if cuda %{ -DSYCL_EXT_ONEAPI_BACKEND_CUDA_EXPERIMENTAL %} %else %{ %if level_zero %{ -DSYCL_EXT_ONEAPI_BACKEND_L0 %} %} %} -o %t.out
+// XFAIL: any-device-is-opencl, any-device-is-cuda, (windows && any-device-is-level_zero), gpu-intel-dg2, hip_amd
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/15819
+// RUN: %if any-device-is-opencl %{ %{build} -o %t-opencl.out %}
+// RUN: %if any-device-is-level_zero %{ %{build} -DBUILD_FOR_L0 -o %t-l0.out %}
+// RUN: %if any-device-is-cuda %{ %{build} -DBUILD_FOR_CUDA -o %t-cuda.out %}
+// RUN: %if any-device-is-hip %{ %{build} -DBUILD_FOR_HIP -o %t-hip.out %}
 
 #include <sycl/backend.hpp>
 #include <sycl/detail/core.hpp>
 #include <sycl/properties/all_properties.hpp>
 #include <sycl/usm.hpp>
 using namespace sycl;
 
-#ifdef SYCL_EXT_ONEAPI_BACKEND_CUDA_EXPERIMENTAL
+#ifdef BUILD_FOR_CUDA
 #include <sycl/ext/oneapi/experimental/backend/cuda.hpp>
 constexpr auto BACKEND = backend::ext_oneapi_cuda;
 using nativeDevice = CUdevice;
 using nativeQueue = CUstream;
 using nativeEvent = CUevent;
-#elif defined(SYCL_EXT_ONEAPI_BACKEND_HIP)
+#elif defined(BUILD_FOR_HIP)
 #include <sycl/ext/oneapi/backend/hip.hpp>
 constexpr auto BACKEND = backend::ext_oneapi_hip;
 using nativeDevice = hipDevice_t;
 using nativeQueue = hipStream_t;
 using nativeEvent = hipEvent_t;
-#elif defined(SYCL_EXT_ONEAPI_BACKEND_L0)
+#elif defined(BUILD_FOR_L0)
 constexpr auto BACKEND = backend::ext_oneapi_level_zero;
 using nativeDevice = ze_device_handle_t;
 using nativeQueue = ze_command_queue_handle_t;

sycl/test-e2e/NewOffloadDriver/diamond_shape.cpp

@@ -229,7 +229,6 @@ environment:
  * **dump_ir**: - compiler can / cannot dump IR;
  * **llvm-spirv** - llvm-spirv tool availability;
  * **llvm-link** - llvm-link tool availability;
- * **fusion**: - Runtime supports kernel fusion;
  * **aspect-\<name\>**: - SYCL aspects supported by a device;
  * **arch-\<name\>** - [SYCL architecture](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_device_architecture.asciidoc) of a device (e.g. `arch-intel_gpu_pvc`, the name matches what you
    can pass into `-fsycl-targets` compiler flag);

sycl/test-e2e/README.md

@@ -0,0 +1,109 @@
+// REQUIRES: cuda
+
+// RUN: %{build} -o %t.out
+// RUN: %t.out
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+
+#include <sycl/ext/oneapi/bindless_images.hpp>
+
+// Uncomment to print additional test information
+// #define VERBOSE_PRINT
+
+class image_kernel;
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+int main() {
+  sycl::device dev;
+  sycl::queue q(dev);
+  auto ctxt = q.get_context();
+
+  constexpr size_t width = 512;
+  std::vector<float> out(width);
+  std::vector<float> expected(width);
+  std::vector<sycl::float3> dataIn(width);
+  float exp = 512;
+  for (int i = 0; i < width; i++) {
+    expected[i] = exp;
+    dataIn[i] = sycl::float3(exp, width, i);
+  }
+
+  try {
+    // Main point of this test is to check creating an image
+    // with a 3-channel format
+    syclexp::image_descriptor desc({width}, 3, sycl::image_channel_type::fp32);
+
+    syclexp::image_mem imgMem(desc, dev, ctxt);
+
+    q.ext_oneapi_copy(dataIn.data(), imgMem.get_handle(), desc);
+    q.wait_and_throw();
+
+    // Some backends don't support 3-channel formats
+    // We still try to create the image,
+    // but we expect it to fail with UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT
+    syclexp::unsampled_image_handle imgHandle =
+        sycl::ext::oneapi::experimental::create_image(imgMem, desc, dev, ctxt);
+
+    sycl::buffer<float> buf(out.data(), width);
+
+    q.submit([&](sycl::handler &cgh) {
+      sycl::accessor outAcc{buf, cgh};
+
+      cgh.parallel_for<image_kernel>(width, [=](sycl::id<1> id) {
+#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
+        // This shouldn't be hit anyway since CUDA doesn't support
+        // 3-channel formats, but we need to ensure the kernel can compile
+        using pixel_t = sycl::float4;
+#else
+        using pixel_t = sycl::float3;
+#endif
+        auto pixel = syclexp::fetch_image<pixel_t>(imgHandle, int(id[0]));
+        outAcc[id] = pixel[0];
+      });
+    });
+    q.wait_and_throw();
+
+  } catch (const sycl::exception &ex) {
+    const std::string_view errMsg(ex.what());
+    if (ctxt.get_backend() == sycl::backend::ext_oneapi_cuda) {
+      if (errMsg.find("UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT") !=
+          std::string::npos) {
+        std::cout << "CUDA doesn't support 3-channel formats, test passed."
+                  << std::endl;
+        return 0;
+      }
+    }
+    std::cerr << "Unexpected SYCL exception: " << errMsg << "\n";
+    return 1;
+  } catch (...) {
+    std::cerr << "Unknown exception caught!\n";
+    return 2;
+  }
+
+  bool validated = true;
+  for (int i = 0; i < width; 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;
+}