[BINDLESS][L0][E2E] set SAMPLED_IMAGE_FETCH_1D_SUPPORT true

sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D.cpp

@@ -0,0 +1,11 @@
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_1d
+// UNSUPPORTED: target-amd
+// UNSUPPORTED-INTENDED: Sampled fetch not currently supported on AMD
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
+
+#include "fetch_1D.hpp"
+
+int main() { return test(); }

sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D.hpp

@@ -0,0 +1,92 @@
+#include <iostream>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/bindless_images.hpp>
+
+class kernel_sampled_fetch;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+int test() {
+
+  sycl::queue q{};
+
+  // declare image data
+  constexpr size_t N = 30;
+  std::vector<float> out(N);
+  std::vector<float> dataIn(N);
+  for (int i = 0; i < N; i++) {
+    dataIn[i] = i;
+  }
+
+  try {
+    syclexp::bindless_image_sampler samp(
+        sycl::addressing_mode::repeat,
+        sycl::coordinate_normalization_mode::unnormalized,
+        sycl::filtering_mode::nearest);
+
+    // Extension: image descriptor
+    syclexp::image_descriptor desc(N, 1, sycl::image_channel_type::fp32);
+
+    // Extension: allocate memory on device
+    syclexp::image_mem imgMem(desc, q);
+
+    // Extension: copy over data to device for non-USM image
+    q.ext_oneapi_copy(dataIn.data(), imgMem.get_handle(), desc);
+    q.wait_and_throw();
+
+    // Extension: create the images and return the handles
+    syclexp::sampled_image_handle imgHandle =
+        syclexp::create_image(imgMem, samp, desc, q);
+
+    sycl::buffer buf(out.data(), sycl::range{N});
+    q.submit([&](sycl::handler &cgh) {
+      auto outAcc =
+          buf.get_access<sycl::access_mode::write>(cgh, sycl::range<1>{N});
+
+      cgh.parallel_for<kernel_sampled_fetch>(
+          sycl::nd_range<1>{N, N}, [=](sycl::nd_item<1> it) {
+            size_t dim0 = it.get_local_id(0);
+            // Extension: fetch data from sampled image handle
+            outAcc[dim0] = syclexp::fetch_image<float>(imgHandle, int(dim0));
+          });
+    });
+
+    q.wait_and_throw();
+
+    // Extension: cleanup
+    syclexp::destroy_image_handle(imgHandle, q);
+  } 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] != dataIn[i]) {
+      mismatch = true;
+      validated = false;
+    }
+
+    if (mismatch) {
+#ifdef VERBOSE_PRINT
+      std::cout << "Result mismatch! Expected: " << dataIn[i]
+                << ", Actual: " << out[i] << "\n";
+#else
+      break;
+#endif
+    }
+  }
+  if (validated) {
+    std::cout << "Test passed!"
+              << "\n";
+    return 0;
+  }
+
+  std::cout << "Test failed!"
+            << "\n";
+  return 3;
+}

sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D_O0.cpp

@@ -0,0 +1,13 @@
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_1d
+// UNSUPPORTED: target-amd
+// UNSUPPORTED-INTENDED: Sampled fetch not currently supported on AMD
+// XFAIL: level_zero && windows
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18919
+
+// RUN: %{build} %O0 -o %t.out
+// RUN: %{run} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
+
+#include "fetch_1D.hpp"
+
+int main() { return test(); }

unified-runtime/source/adapters/level_zero/device.cpp

@@ -1149,8 +1149,8 @@ ur_result_t urDeviceGetInfo(
     return ReturnValue(true);
   }
   case UR_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_1D_SUPPORT_EXP: {
-    // L0 does not support fetching 1D non-USM sampled image data.
-    return ReturnValue(false);
+    // L0 does support fetching 1D non-USM sampled image data.
+    return ReturnValue(true);
   }
   case UR_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_2D_USM_SUPPORT_EXP: {
     // L0 does support fetching 2D USM sampled image data.