From 5f0230dab198b991f95f05f847b848b4baffddb8 Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Thu, 15 May 2025 12:22:09 +0100 Subject: [PATCH 1/4] Add l0 failing O0 bindless tests Signed-off-by: JackAKirk --- sycl/test-e2e/bindless_images/read_1D.cpp | 142 +----------------- sycl/test-e2e/bindless_images/read_1D.hpp | 141 +++++++++++++++++ sycl/test-e2e/bindless_images/read_1D_O0.cpp | 13 ++ .../sampled_fetch/fetch_2D.cpp | 108 +------------ .../sampled_fetch/fetch_2D.hpp | 107 +++++++++++++ .../sampled_fetch/fetch_2D_O0.cpp | 12 ++ sycl/test-e2e/bindless_images/sampling_1D.cpp | 110 +------------- sycl/test-e2e/bindless_images/sampling_1D.hpp | 109 ++++++++++++++ .../bindless_images/sampling_1D_O0.cpp | 15 ++ 9 files changed, 403 insertions(+), 354 deletions(-) create mode 100644 sycl/test-e2e/bindless_images/read_1D.hpp create mode 100644 sycl/test-e2e/bindless_images/read_1D_O0.cpp create mode 100644 sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D.hpp create mode 100644 sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_O0.cpp create mode 100644 sycl/test-e2e/bindless_images/sampling_1D.hpp create mode 100644 sycl/test-e2e/bindless_images/sampling_1D_O0.cpp 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..d4ea41cca596e --- /dev/null +++ b/sycl/test-e2e/bindless_images/read_1D_O0.cpp @@ -0,0 +1,13 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images +// XFAIL: level_zero + +// 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..a43bf351e8c80 --- /dev/null +++ b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_O0.cpp @@ -0,0 +1,12 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images +// REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_2d +// XFAIL: level_zero +// 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..6c8562df17059 --- /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 + +// 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(); } From 2aeb5b38d6b639fc9988e9f7356c8ef5bc4a6b7e Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Mon, 19 May 2025 15:02:19 +0100 Subject: [PATCH 2/4] Small reproducer CI check Test reproducer on latest l0 CI drivers Signed-off-by: JackAKirk --- .../examples/ptrptr_handle.cpp | 97 +++++++++++++++++++ 1 file changed, 97 insertions(+) create mode 100644 sycl/test-e2e/bindless_images/examples/ptrptr_handle.cpp diff --git a/sycl/test-e2e/bindless_images/examples/ptrptr_handle.cpp b/sycl/test-e2e/bindless_images/examples/ptrptr_handle.cpp new file mode 100644 index 0000000000000..a6bc517c6eb7a --- /dev/null +++ b/sycl/test-e2e/bindless_images/examples/ptrptr_handle.cpp @@ -0,0 +1,97 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images +// UNSUPPORTED: hip +// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test. + +// RUN: %{build} -o %t.out +// RUN: %{run} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out + +#include +#include +#include + +int main() { + // Set up device, queue, and context + sycl::device dev; + sycl::queue q(dev); + sycl::context ctxt = q.get_context(); + + // Initialize input data + constexpr size_t width = 512; + std::vector dataIn(width); + std::vector dataOut(width); + for (int i = 0; i < width; i++) { + dataIn[i] = static_cast(i); + } + + // Image descriptor - can use the same for both images + sycl::ext::oneapi::experimental::image_descriptor desc( + sycl::range{width}, 1, sycl::image_channel_type::fp32); + + // Extension: returns the device pointer to the allocated memory + sycl::ext::oneapi::experimental::image_mem imgMemoryIn(desc, q); + sycl::ext::oneapi::experimental::image_mem imgMemoryOut(desc, q); + + // Extension: create the image and return the handle + sycl::ext::oneapi::experimental::unsampled_image_handle imgIn = + sycl::ext::oneapi::experimental::create_image(imgMemoryIn, desc, q); + sycl::ext::oneapi::experimental::unsampled_image_handle imgOut = + sycl::ext::oneapi::experimental::create_image(imgMemoryOut, desc, q); + +void * Texptr=reinterpret_cast(sycl::malloc_device(sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle), q)); +q.memcpy(reinterpret_cast(Texptr), + reinterpret_cast(&imgIn), sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle)); + +q.ext_oneapi_copy(dataIn.data(), imgMemoryIn.get_handle(), desc); +q.wait(); + +void * Texptr_ptr=reinterpret_cast(sycl::malloc_device(sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle*), q)); +q.memcpy(reinterpret_cast(Texptr_ptr), + &Texptr, sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle*)); +q.wait(); + + + // Bindless images require manual synchronization + // Wait for copy operation to finish + q.wait_and_throw(); + + q.submit([&](sycl::handler &cgh) { + // No need to request access, handles captured by value + + cgh.parallel_for( + sycl::nd_range<1>{{width}, {width}}, [=](sycl::nd_item<1> it) { + +sycl::ext::oneapi::experimental::unsampled_image_handle** image_array = reinterpret_cast(Texptr_ptr); +sycl::ext::oneapi::experimental::unsampled_image_handle* imgIn2p=reinterpret_cast(image_array[0]); +sycl::ext::oneapi::experimental::unsampled_image_handle imgIn2=imgIn2p[0]; + size_t dim0 = it.get_local_id(0); + // Extension: read image data from handle + float pixel = sycl::ext::oneapi::experimental::fetch_image( + imgIn2, int(dim0)); + + // Extension: write to image data using handle + sycl::ext::oneapi::experimental::write_image(imgOut, int(dim0), + pixel); + }); + }); + + // Using image handles requires manual synchronization + q.wait_and_throw(); + + // Copy data written to imgOut to host + q.ext_oneapi_copy(imgMemoryOut.get_handle(), dataOut.data(), desc); + + // Ensure copying data from the device to host is finished before validate + q.wait_and_throw(); + + // Cleanup + sycl::ext::oneapi::experimental::destroy_image_handle(imgIn, q); + sycl::ext::oneapi::experimental::destroy_image_handle(imgOut, q); + + for (size_t i = 0; i < width; i++) { +//std::cout << dataOut[i] << dataIn[i] << "\n"; + if (dataOut[i] != dataIn[i]) { + return 1; + } + } + return 0; +} From 9419e6a3543ce2ce5b772353da50c1b158abbac3 Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Wed, 11 Jun 2025 11:11:12 +0100 Subject: [PATCH 3/4] xfail windows Signed-off-by: JackAKirk --- .../examples/ptrptr_handle.cpp | 97 ------------------- sycl/test-e2e/bindless_images/read_1D_O0.cpp | 2 +- .../sampled_fetch/fetch_2D_O0.cpp | 2 +- .../bindless_images/sampling_1D_O0.cpp | 2 +- 4 files changed, 3 insertions(+), 100 deletions(-) delete mode 100644 sycl/test-e2e/bindless_images/examples/ptrptr_handle.cpp diff --git a/sycl/test-e2e/bindless_images/examples/ptrptr_handle.cpp b/sycl/test-e2e/bindless_images/examples/ptrptr_handle.cpp deleted file mode 100644 index a6bc517c6eb7a..0000000000000 --- a/sycl/test-e2e/bindless_images/examples/ptrptr_handle.cpp +++ /dev/null @@ -1,97 +0,0 @@ -// REQUIRES: aspect-ext_oneapi_bindless_images -// UNSUPPORTED: hip -// UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test. - -// RUN: %{build} -o %t.out -// RUN: %{run} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out - -#include -#include -#include - -int main() { - // Set up device, queue, and context - sycl::device dev; - sycl::queue q(dev); - sycl::context ctxt = q.get_context(); - - // Initialize input data - constexpr size_t width = 512; - std::vector dataIn(width); - std::vector dataOut(width); - for (int i = 0; i < width; i++) { - dataIn[i] = static_cast(i); - } - - // Image descriptor - can use the same for both images - sycl::ext::oneapi::experimental::image_descriptor desc( - sycl::range{width}, 1, sycl::image_channel_type::fp32); - - // Extension: returns the device pointer to the allocated memory - sycl::ext::oneapi::experimental::image_mem imgMemoryIn(desc, q); - sycl::ext::oneapi::experimental::image_mem imgMemoryOut(desc, q); - - // Extension: create the image and return the handle - sycl::ext::oneapi::experimental::unsampled_image_handle imgIn = - sycl::ext::oneapi::experimental::create_image(imgMemoryIn, desc, q); - sycl::ext::oneapi::experimental::unsampled_image_handle imgOut = - sycl::ext::oneapi::experimental::create_image(imgMemoryOut, desc, q); - -void * Texptr=reinterpret_cast(sycl::malloc_device(sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle), q)); -q.memcpy(reinterpret_cast(Texptr), - reinterpret_cast(&imgIn), sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle)); - -q.ext_oneapi_copy(dataIn.data(), imgMemoryIn.get_handle(), desc); -q.wait(); - -void * Texptr_ptr=reinterpret_cast(sycl::malloc_device(sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle*), q)); -q.memcpy(reinterpret_cast(Texptr_ptr), - &Texptr, sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle*)); -q.wait(); - - - // Bindless images require manual synchronization - // Wait for copy operation to finish - q.wait_and_throw(); - - q.submit([&](sycl::handler &cgh) { - // No need to request access, handles captured by value - - cgh.parallel_for( - sycl::nd_range<1>{{width}, {width}}, [=](sycl::nd_item<1> it) { - -sycl::ext::oneapi::experimental::unsampled_image_handle** image_array = reinterpret_cast(Texptr_ptr); -sycl::ext::oneapi::experimental::unsampled_image_handle* imgIn2p=reinterpret_cast(image_array[0]); -sycl::ext::oneapi::experimental::unsampled_image_handle imgIn2=imgIn2p[0]; - size_t dim0 = it.get_local_id(0); - // Extension: read image data from handle - float pixel = sycl::ext::oneapi::experimental::fetch_image( - imgIn2, int(dim0)); - - // Extension: write to image data using handle - sycl::ext::oneapi::experimental::write_image(imgOut, int(dim0), - pixel); - }); - }); - - // Using image handles requires manual synchronization - q.wait_and_throw(); - - // Copy data written to imgOut to host - q.ext_oneapi_copy(imgMemoryOut.get_handle(), dataOut.data(), desc); - - // Ensure copying data from the device to host is finished before validate - q.wait_and_throw(); - - // Cleanup - sycl::ext::oneapi::experimental::destroy_image_handle(imgIn, q); - sycl::ext::oneapi::experimental::destroy_image_handle(imgOut, q); - - for (size_t i = 0; i < width; i++) { -//std::cout << dataOut[i] << dataIn[i] << "\n"; - if (dataOut[i] != dataIn[i]) { - return 1; - } - } - return 0; -} diff --git a/sycl/test-e2e/bindless_images/read_1D_O0.cpp b/sycl/test-e2e/bindless_images/read_1D_O0.cpp index d4ea41cca596e..1e5e0e09cfae3 100644 --- a/sycl/test-e2e/bindless_images/read_1D_O0.cpp +++ b/sycl/test-e2e/bindless_images/read_1D_O0.cpp @@ -1,5 +1,5 @@ // REQUIRES: aspect-ext_oneapi_bindless_images -// XFAIL: level_zero +// XFAIL: level_zero && windows // RUN: %{build} %O0 -o %t.out // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out 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 index a43bf351e8c80..f4167b8513a33 100644 --- a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_O0.cpp +++ b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_O0.cpp @@ -1,6 +1,6 @@ // REQUIRES: aspect-ext_oneapi_bindless_images // REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_2d -// XFAIL: level_zero +// XFAIL: level_zero && windows // UNSUPPORTED: target-amd // UNSUPPORTED-INTENDED: Sampled fetch not currently supported on AMD diff --git a/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp b/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp index 6c8562df17059..80774b7151abc 100644 --- a/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp +++ b/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp @@ -1,5 +1,5 @@ // REQUIRES: aspect-ext_oneapi_bindless_images -// XFAIL: level_zero +// XFAIL: level_zero && windows // RUN: %{build} %O0 -o %t.out // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out From 42b108268e8e98eeb1227b3c7041d01f716ee765 Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Wed, 11 Jun 2025 13:51:57 +0100 Subject: [PATCH 4/4] Add XFAIL trackers Signed-off-by: JackAKirk --- sycl/test-e2e/bindless_images/read_1D_O0.cpp | 3 +-- sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_O0.cpp | 1 + sycl/test-e2e/bindless_images/sampling_1D_O0.cpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/bindless_images/read_1D_O0.cpp b/sycl/test-e2e/bindless_images/read_1D_O0.cpp index 1e5e0e09cfae3..d25911de413ce 100644 --- a/sycl/test-e2e/bindless_images/read_1D_O0.cpp +++ b/sycl/test-e2e/bindless_images/read_1D_O0.cpp @@ -1,6 +1,6 @@ // 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 @@ -10,4 +10,3 @@ #include "read_1D.hpp" int main() { return test(); } - 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 index f4167b8513a33..11647caaa919b 100644 --- a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_O0.cpp +++ b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_O0.cpp @@ -1,6 +1,7 @@ // 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 diff --git a/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp b/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp index 80774b7151abc..b96dd9e7c8af8 100644 --- a/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp +++ b/sycl/test-e2e/bindless_images/sampling_1D_O0.cpp @@ -1,6 +1,6 @@ // 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