diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D.cpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D.cpp new file mode 100644 index 0000000000000..a1970bc8c634c --- /dev/null +++ b/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(); } diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D.hpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D.hpp new file mode 100644 index 0000000000000..c104f9291a237 --- /dev/null +++ b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D.hpp @@ -0,0 +1,92 @@ +#include +#include +#include + +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 out(N); + std::vector 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(cgh, sycl::range<1>{N}); + + cgh.parallel_for( + 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(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; +} diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D_O0.cpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D_O0.cpp new file mode 100644 index 0000000000000..2a25d5f0a6956 --- /dev/null +++ b/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(); } diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index 5dd107a47608c..05e808a5aabcc 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/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.