diff --git a/sycl/test-e2e/bindless_images/read_1D.cpp b/sycl/test-e2e/bindless_images/read_1D.cpp index b619710fa2f1a..0d6974441199f 100644 --- a/sycl/test-e2e/bindless_images/read_1D.cpp +++ b/sycl/test-e2e/bindless_images/read_1D.cpp @@ -3,147 +3,9 @@ // RUN: %{build} -o %t.out // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out -#include -#include - -#include - // 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 width = 512; - std::vector out(width); - std::vector expected(width); - std::vector dataIn1(width); - std::vector dataIn2(width); - float exp = 512; - for (int i = 0; i < width; i++) { - expected[i] = exp; - dataIn1[i] = sycl::float4(i, i, i, i); - dataIn2[i] = sycl::float4(width - i, width - i, width - i, width - i); - } - - try { - // Extension: image descriptor - can use the same for both images - sycl::ext::oneapi::experimental::image_descriptor desc( - {width}, 4, sycl::image_channel_type::fp32); - - // Extension: allocate memory on device and create the handle - sycl::ext::oneapi::experimental::image_mem imgMem0(desc, dev, ctxt); - sycl::ext::oneapi::experimental::image_mem imgMem1(desc, dev, ctxt); - - // std::hash specialization to ensure `image_mem` follows common reference - // semantics - assert(std::hash{}(imgMem0) != - std::hash{}(imgMem1)); - - // We're able to use move semantics - // Move construct - sycl::ext::oneapi::experimental::image_mem imgMem0MoveConstruct( - std::move(imgMem0)); - // Move assign - sycl::ext::oneapi::experimental::image_mem imgMem0MoveAssign; - imgMem0MoveAssign = std::move(imgMem0MoveConstruct); - - // We're able to use copy semantics - // Copy construct - sycl::ext::oneapi::experimental::image_mem imgMem1CopyConstruct(imgMem1); - // Copy assign - sycl::ext::oneapi::experimental::image_mem imgMem1CopyAssign; - imgMem1CopyAssign = imgMem1CopyConstruct; - - // Equality operators to ensure `image_mem` follows common reference - // semantics - assert(imgMem0MoveAssign != imgMem1CopyAssign); - assert(imgMem1 == imgMem1CopyAssign); - - // We can default construct image handles - sycl::ext::oneapi::experimental::unsampled_image_handle imgHandle1; - - // Extension: create the image and return the handle - sycl::ext::oneapi::experimental::unsampled_image_handle tmpHandle = - sycl::ext::oneapi::experimental::create_image(imgMem0MoveAssign, desc, - dev, ctxt); - sycl::ext::oneapi::experimental::unsampled_image_handle imgHandle2 = - sycl::ext::oneapi::experimental::create_image(imgMem1CopyAssign, desc, - dev, ctxt); - - // Default constructed image handles are not valid until we assign a valid - // raw handle to the struct - imgHandle1.raw_handle = tmpHandle.raw_handle; - - // Extension: copy over data to device - q.ext_oneapi_copy(dataIn1.data(), imgMem0MoveAssign.get_handle(), desc); - q.ext_oneapi_copy(dataIn2.data(), imgMem1CopyAssign.get_handle(), desc); - - q.wait_and_throw(); - - sycl::buffer buf((float *)out.data(), width); - q.submit([&](sycl::handler &cgh) { - auto outAcc = buf.get_access(cgh, width); - - cgh.parallel_for(width, [=](sycl::id<1> id) { - float sum = 0; - // Extension: fetch image data from handle - sycl::float4 px1 = - sycl::ext::oneapi::experimental::fetch_image( - imgHandle1, int(id[0])); - sycl::float4 px2 = - sycl::ext::oneapi::experimental::fetch_image( - imgHandle2, int(id[0])); - - sum = px1[0] + px2[0]; - outAcc[id] = sum; - }); - }); - - q.wait_and_throw(); - - // Extension: cleanup - sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle1, dev, - ctxt); - sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle2, 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 < 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; - } +#include "read_1D.hpp" - std::cout << "Test failed!" << std::endl; - return 3; -} +int main() { return test(); } diff --git a/sycl/test-e2e/bindless_images/read_1D.hpp b/sycl/test-e2e/bindless_images/read_1D.hpp new file mode 100644 index 0000000000000..14b6ae722e2ac --- /dev/null +++ b/sycl/test-e2e/bindless_images/read_1D.hpp @@ -0,0 +1,141 @@ +#include +#include + +#include + +class image_addition; + +int test() { + + sycl::device dev; + sycl::queue q(dev); + auto ctxt = q.get_context(); + + // declare image data + constexpr size_t width = 512; + std::vector out(width); + std::vector expected(width); + std::vector dataIn1(width); + std::vector dataIn2(width); + float exp = 512; + for (int i = 0; i < width; i++) { + expected[i] = exp; + dataIn1[i] = sycl::float4(i, i, i, i); + dataIn2[i] = sycl::float4(width - i, width - i, width - i, width - i); + } + + try { + // Extension: image descriptor - can use the same for both images + sycl::ext::oneapi::experimental::image_descriptor desc( + {width}, 4, sycl::image_channel_type::fp32); + + // Extension: allocate memory on device and create the handle + sycl::ext::oneapi::experimental::image_mem imgMem0(desc, dev, ctxt); + sycl::ext::oneapi::experimental::image_mem imgMem1(desc, dev, ctxt); + + // std::hash specialization to ensure `image_mem` follows common reference + // semantics + assert(std::hash{}(imgMem0) != + std::hash{}(imgMem1)); + + // We're able to use move semantics + // Move construct + sycl::ext::oneapi::experimental::image_mem imgMem0MoveConstruct( + std::move(imgMem0)); + // Move assign + sycl::ext::oneapi::experimental::image_mem imgMem0MoveAssign; + imgMem0MoveAssign = std::move(imgMem0MoveConstruct); + + // We're able to use copy semantics + // Copy construct + sycl::ext::oneapi::experimental::image_mem imgMem1CopyConstruct(imgMem1); + // Copy assign + sycl::ext::oneapi::experimental::image_mem imgMem1CopyAssign; + imgMem1CopyAssign = imgMem1CopyConstruct; + + // Equality operators to ensure `image_mem` follows common reference + // semantics + assert(imgMem0MoveAssign != imgMem1CopyAssign); + assert(imgMem1 == imgMem1CopyAssign); + + // We can default construct image handles + sycl::ext::oneapi::experimental::unsampled_image_handle imgHandle1; + + // Extension: create the image and return the handle + sycl::ext::oneapi::experimental::unsampled_image_handle tmpHandle = + sycl::ext::oneapi::experimental::create_image(imgMem0MoveAssign, desc, + dev, ctxt); + sycl::ext::oneapi::experimental::unsampled_image_handle imgHandle2 = + sycl::ext::oneapi::experimental::create_image(imgMem1CopyAssign, desc, + dev, ctxt); + + // Default constructed image handles are not valid until we assign a valid + // raw handle to the struct + imgHandle1.raw_handle = tmpHandle.raw_handle; + + // Extension: copy over data to device + q.ext_oneapi_copy(dataIn1.data(), imgMem0MoveAssign.get_handle(), desc); + q.ext_oneapi_copy(dataIn2.data(), imgMem1CopyAssign.get_handle(), desc); + + q.wait_and_throw(); + + sycl::buffer buf((float *)out.data(), width); + q.submit([&](sycl::handler &cgh) { + auto outAcc = buf.get_access(cgh, width); + + cgh.parallel_for(width, [=](sycl::id<1> id) { + float sum = 0; + // Extension: fetch image data from handle + sycl::float4 px1 = + sycl::ext::oneapi::experimental::fetch_image( + imgHandle1, int(id[0])); + sycl::float4 px2 = + sycl::ext::oneapi::experimental::fetch_image( + imgHandle2, int(id[0])); + + sum = px1[0] + px2[0]; + outAcc[id] = sum; + }); + }); + + q.wait_and_throw(); + + // Extension: cleanup + sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle1, dev, + ctxt); + sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle2, 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 < 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; +} diff --git a/sycl/test-e2e/bindless_images/read_1D_O0.cpp b/sycl/test-e2e/bindless_images/read_1D_O0.cpp new file mode 100644 index 0000000000000..d25911de413ce --- /dev/null +++ b/sycl/test-e2e/bindless_images/read_1D_O0.cpp @@ -0,0 +1,12 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images +// XFAIL: level_zero && windows +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18919 +// RUN: %{build} %O0 -o %t.out +// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out + +// Uncomment to print additional test information +// #define VERBOSE_PRINT + +#include "read_1D.hpp" + +int main() { return test(); } diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.cpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.cpp index e1970128f2fe9..cdd9e1fc0521b 100644 --- a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.cpp +++ b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.cpp @@ -6,110 +6,6 @@ // RUN: %{build} -o %t.out // RUN: %{run} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out -#include -#include -#include +#include "fetch_2D.hpp" -class kernel_sampled_fetch; - -int main() { - - sycl::device dev; - sycl::queue q(dev); - auto ctxt = q.get_context(); - - // declare image data - constexpr size_t width = 5; - constexpr size_t height = 6; - constexpr size_t N = width * height; - std::vector out(N); - std::vector expected(N); - std::vector dataIn(N); - for (int i = 0; i < width; i++) { - for (int j = 0; j < height; j++) { - auto index = i + (width * j); - expected[index] = index; - dataIn[index] = index; - } - } - - namespace syclexp = sycl::ext::oneapi::experimental; - - 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({width, height}, 1, - sycl::image_channel_type::fp32); - - // Extension: allocate memory on device - syclexp::image_mem imgMem(desc, dev, ctxt); - - // 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{height, width}); - q.submit([&](sycl::handler &cgh) { - auto outAcc = buf.get_access( - cgh, sycl::range<2>{height, width}); - - cgh.parallel_for( - sycl::nd_range<2>{{width, height}, {width, height}}, - [=](sycl::nd_item<2> it) { - size_t dim0 = it.get_local_id(0); - size_t dim1 = it.get_local_id(1); - - // Extension: fetch data from sampled image handle - float px1 = - syclexp::fetch_image(imgHandle, sycl::int2(dim0, dim1)); - - outAcc[sycl::id<2>{dim1, dim0}] = px1; - }); - }); - - q.wait_and_throw(); - - // Extension: cleanup - syclexp::destroy_image_handle(imgHandle, 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; -} +int main() { return test(); } diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.hpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.hpp new file mode 100644 index 0000000000000..550c5f877d256 --- /dev/null +++ b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.hpp @@ -0,0 +1,107 @@ +#include +#include +#include + +class kernel_sampled_fetch; + +int test() { + + sycl::device dev; + sycl::queue q(dev); + auto ctxt = q.get_context(); + + // declare image data + constexpr size_t width = 5; + constexpr size_t height = 6; + constexpr size_t N = width * height; + std::vector out(N); + std::vector expected(N); + std::vector dataIn(N); + for (int i = 0; i < width; i++) { + for (int j = 0; j < height; j++) { + auto index = i + (width * j); + expected[index] = index; + dataIn[index] = index; + } + } + + namespace syclexp = sycl::ext::oneapi::experimental; + + 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({width, height}, 1, + sycl::image_channel_type::fp32); + + // Extension: allocate memory on device + syclexp::image_mem imgMem(desc, dev, ctxt); + + // 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{height, width}); + q.submit([&](sycl::handler &cgh) { + auto outAcc = buf.get_access( + cgh, sycl::range<2>{height, width}); + + cgh.parallel_for( + sycl::nd_range<2>{{width, height}, {width, height}}, + [=](sycl::nd_item<2> it) { + size_t dim0 = it.get_local_id(0); + size_t dim1 = it.get_local_id(1); + + // Extension: fetch data from sampled image handle + float px1 = + syclexp::fetch_image(imgHandle, sycl::int2(dim0, dim1)); + + outAcc[sycl::id<2>{dim1, dim0}] = px1; + }); + }); + + q.wait_and_throw(); + + // Extension: cleanup + syclexp::destroy_image_handle(imgHandle, 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; +} diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_O0.cpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_O0.cpp new file mode 100644 index 0000000000000..11647caaa919b --- /dev/null +++ b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_O0.cpp @@ -0,0 +1,13 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images +// REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_2d +// XFAIL: level_zero && windows +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18919 +// UNSUPPORTED: target-amd +// UNSUPPORTED-INTENDED: Sampled fetch not currently supported on AMD + +// RUN: %{build} %O0 -o %t.out +// RUN: %{run} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out + +#include "fetch_2D.hpp" + +int main() { return test(); } diff --git a/sycl/test-e2e/bindless_images/sampling_1D.cpp b/sycl/test-e2e/bindless_images/sampling_1D.cpp index ea9f667bdb372..ab99babc0d7b6 100644 --- a/sycl/test-e2e/bindless_images/sampling_1D.cpp +++ b/sycl/test-e2e/bindless_images/sampling_1D.cpp @@ -6,115 +6,9 @@ // UNSUPPORTED: hip // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17212 -#include -#include -#include - // Uncomment to print additional test information // #define VERBOSE_PRINT -class image_addition; - -int main() { - -#if defined(SYCL_EXT_ONEAPI_BINDLESS_IMAGES) - assert(SYCL_EXT_ONEAPI_BINDLESS_IMAGES == 6); -#if defined(VERBOSE_PRINT) - std::cout << "SYCL_EXT_ONEAPI_BINDLESS_IMAGES is defined!" << std::endl; -#endif -#else - std::cerr << "Bindless images feature test macro is not defined!" - << std::endl; - return 1; -#endif // defined(SYCL_EXT_ONEAPI_BINDLESS_IMAGES) - - sycl::device dev; - sycl::queue q(dev); - auto ctxt = q.get_context(); - - // declare image data - constexpr size_t N = 32; - size_t width = N; - std::vector out(N); - std::vector expected(N); - std::vector dataIn(N); - for (int i = 0; i < N; i++) { - expected[i] = i; - dataIn[i] = float(i); - } - - try { - // Extension: image descriptor - sycl::ext::oneapi::experimental::image_descriptor desc( - {width}, 1, sycl::image_channel_type::fp32); - - sycl::ext::oneapi::experimental::bindless_image_sampler samp( - sycl::addressing_mode::repeat, - sycl::coordinate_normalization_mode::normalized, - sycl::filtering_mode::linear); - - // Extension: allocate memory on device - sycl::ext::oneapi::experimental::image_mem imgMem(desc, dev, ctxt); - - // Extension: copy over data to device - q.ext_oneapi_copy(dataIn.data(), imgMem.get_handle(), desc); - q.wait_and_throw(); - - // Extension: create the image and return the handle - auto imgHandle = sycl::ext::oneapi::experimental::create_image( - imgMem, samp, desc, dev, ctxt); - - sycl::buffer buf((float *)out.data(), N); - q.submit([&](sycl::handler &cgh) { - auto outAcc = buf.get_access(cgh, N); - - cgh.parallel_for(N, [=](sycl::id<1> id) { - // Normalize coordinate -- +0.5 to look towards centre of pixel - float x = float(id[0] + 0.5f) / (float)N; - // Extension: sample image data from handle - float px1 = - sycl::ext::oneapi::experimental::sample_image(imgHandle, x); - - outAcc[id] = px1; - }); - }); - - q.wait_and_throw(); - - // Extension: cleanup - sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle, 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; - } +#include "sampling_1D.hpp" - std::cout << "Test failed!" << std::endl; - return 3; -} +int main() { return test(); } diff --git a/sycl/test-e2e/bindless_images/sampling_1D.hpp b/sycl/test-e2e/bindless_images/sampling_1D.hpp new file mode 100644 index 0000000000000..4769f14182bbf --- /dev/null +++ b/sycl/test-e2e/bindless_images/sampling_1D.hpp @@ -0,0 +1,109 @@ +#include +#include +#include + +class image_addition; + +int test() { + +#if defined(SYCL_EXT_ONEAPI_BINDLESS_IMAGES) + assert(SYCL_EXT_ONEAPI_BINDLESS_IMAGES == 6); +#if defined(VERBOSE_PRINT) + std::cout << "SYCL_EXT_ONEAPI_BINDLESS_IMAGES is defined!" << std::endl; +#endif +#else + std::cerr << "Bindless images feature test macro is not defined!" + << std::endl; + return 1; +#endif // defined(SYCL_EXT_ONEAPI_BINDLESS_IMAGES) + + sycl::device dev; + sycl::queue q(dev); + auto ctxt = q.get_context(); + + // declare image data + constexpr size_t N = 32; + size_t width = N; + std::vector out(N); + std::vector expected(N); + std::vector dataIn(N); + for (int i = 0; i < N; i++) { + expected[i] = i; + dataIn[i] = float(i); + } + + try { + // Extension: image descriptor + sycl::ext::oneapi::experimental::image_descriptor desc( + {width}, 1, sycl::image_channel_type::fp32); + + sycl::ext::oneapi::experimental::bindless_image_sampler samp( + sycl::addressing_mode::repeat, + sycl::coordinate_normalization_mode::normalized, + sycl::filtering_mode::linear); + + // Extension: allocate memory on device + sycl::ext::oneapi::experimental::image_mem imgMem(desc, dev, ctxt); + + // Extension: copy over data to device + q.ext_oneapi_copy(dataIn.data(), imgMem.get_handle(), desc); + q.wait_and_throw(); + + // Extension: create the image and return the handle + auto imgHandle = sycl::ext::oneapi::experimental::create_image( + imgMem, samp, desc, dev, ctxt); + + sycl::buffer buf((float *)out.data(), N); + q.submit([&](sycl::handler &cgh) { + auto outAcc = buf.get_access(cgh, N); + + cgh.parallel_for(N, [=](sycl::id<1> id) { + // Normalize coordinate -- +0.5 to look towards centre of pixel + float x = float(id[0] + 0.5f) / (float)N; + // Extension: sample image data from handle + float px1 = + sycl::ext::oneapi::experimental::sample_image(imgHandle, x); + + outAcc[id] = px1; + }); + }); + + q.wait_and_throw(); + + // Extension: cleanup + sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle, 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; +} diff --git a/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp b/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp new file mode 100644 index 0000000000000..b96dd9e7c8af8 --- /dev/null +++ b/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp @@ -0,0 +1,15 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images +// XFAIL: level_zero && windows +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18919 +// RUN: %{build} %O0 -o %t.out +// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out + +// UNSUPPORTED: hip +// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17212 + +// Uncomment to print additional test information +// #define VERBOSE_PRINT + +#include "sampling_1D.hpp" + +int main() { return test(); }