From f40184d86da6712bab63c560ec03ef5aada55562 Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Thu, 29 May 2025 10:11:23 +0100 Subject: [PATCH 1/7] Add ptr to handle ptr test failing on l0 Signed-off-by: JackAKirk --- .../array/fetch_handle_carray2d.cpp | 98 +++++++++++++++++++ 1 file changed, 98 insertions(+) create mode 100644 sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp diff --git a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp new file mode 100644 index 0000000000000..97a07f6291360 --- /dev/null +++ b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp @@ -0,0 +1,98 @@ +// REQUIRES: aspect-ext_oneapi_bindless_images +// XFAIL: level_zero +// 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 + +int main() { + + sycl::queue q{}; + + // 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 *imageHandlePtrGen = static_cast(sycl::malloc_device( + sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle), q)); + q.memcpy(static_cast(imageHandlePtrGen), + static_cast(&imgIn), + sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle)); + q.wait_and_throw(); + q.ext_oneapi_copy(dataIn.data(), imgMemoryIn.get_handle(), desc); + q.wait_and_throw(); + void *imageHandlePtrPtrGen = static_cast(sycl::malloc_device( + sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle *), q)); + q.memcpy(static_cast(imageHandlePtrPtrGen), + static_cast(&imageHandlePtrGen), + sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle *)); + q.wait_and_throw(); + + q.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::nd_range<1>{{width}, {width}}, [=](sycl::nd_item<1> + it) { + sycl::ext::oneapi::experimental::unsampled_image_handle * + *imageHandlePtrPtr = static_cast< + sycl::ext::oneapi::experimental::unsampled_image_handle **>( + imageHandlePtrPtrGen); + sycl::ext::oneapi::experimental::unsampled_image_handle *imageHandlePtr = + static_cast< + sycl::ext::oneapi::experimental::unsampled_image_handle *>( + imageHandlePtrPtr[0]); + sycl::ext::oneapi::experimental::unsampled_image_handle imageHandle = + imageHandlePtr[0]; + + size_t dim0 = it.get_local_id(0); + // Extension: read image data from handle + float pixel = sycl::ext::oneapi::experimental::fetch_image( + imageHandle, int(dim0)); + + // Extension: write to image data using handle + sycl::ext::oneapi::experimental::write_image(imgOut, int(dim0), pixel); + }); + }); + + 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); + sycl::free(imageHandlePtrGen, q); + sycl::free(imageHandlePtrPtrGen, q); + + for (size_t i = 0; i < width; i++) { + if (dataOut[i] != dataIn[i]) { + std::cout << "Test failed" + << "\n"; + return 1; + } + } +return 0; +} From a4e2aae4753816a9bfc074e5d1e09e73cd67d836 Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Thu, 29 May 2025 10:20:51 +0100 Subject: [PATCH 2/7] Fix format Signed-off-by: JackAKirk --- sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp index 97a07f6291360..792df90c7afba 100644 --- a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp +++ b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp @@ -94,5 +94,5 @@ int main() { return 1; } } -return 0; + return 0; } From 7d07c29dd3bc9f40662357ce40380005e18d55d0 Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Thu, 29 May 2025 15:17:20 +0100 Subject: [PATCH 3/7] Fix CI failures, address review comments Signed-off-by: JackAKirk --- .../array/fetch_handle_carray2d.cpp | 101 +++++++++++------- 1 file changed, 61 insertions(+), 40 deletions(-) diff --git a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp index 792df90c7afba..633d9de503a01 100644 --- a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp +++ b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp @@ -1,12 +1,17 @@ // REQUIRES: aspect-ext_oneapi_bindless_images // XFAIL: level_zero +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18727 // 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 +#include + +namespace syclexp = sycl::ext::oneapi::experimental; int main() { @@ -21,56 +26,72 @@ int main() { } // 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); + syclexp::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); + syclexp::image_mem imgMemoryIn(desc, q); + syclexp::image_mem imgMemoryOut(desc, q); + q.ext_oneapi_copy(dataIn.data(), imgMemoryIn.get_handle(), desc); + q.wait_and_throw(); // 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); + syclexp::unsampled_image_handle imgIn = + syclexp::create_image(imgMemoryIn, desc, q); + syclexp::unsampled_image_handle imgOut = + syclexp::create_image(imgMemoryOut, desc, q); + + // Copy the input data to the image_mem of the device unsampled_image_handle + q.ext_oneapi_copy(dataIn.data(), imgMemoryIn.get_handle(), desc); + q.wait_and_throw(); + + // Allocate an unsampled_image_handle manually instead of using create_image + // The implied purpose of this appears to be that a valid device void* is + // returned which points to the image_handle + void *imageHandlePtrGen = + sycl::malloc_device(sizeof(syclexp::unsampled_image_handle), q); - void *imageHandlePtrGen = static_cast(sycl::malloc_device( - sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle), q)); + // Copy the create_image returned device unsampled_image_handle to the contents + // of the void* pointing to the manually created unsampled_image_handle q.memcpy(static_cast(imageHandlePtrGen), static_cast(&imgIn), - sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle)); - q.wait_and_throw(); - q.ext_oneapi_copy(dataIn.data(), imgMemoryIn.get_handle(), desc); + sizeof(syclexp::unsampled_image_handle)); + q.wait_and_throw(); - void *imageHandlePtrPtrGen = static_cast(sycl::malloc_device( - sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle *), q)); + + // Allocate a device generic pointer pointing to an unsampled_image_handle* + void *imageHandlePtrPtrGen = + sycl::malloc_device(sizeof(syclexp::unsampled_image_handle *), q); + + // Copy the address of the manually allocated unsampled_image_handle to the + // contents of the generic device pointer allocated above q.memcpy(static_cast(imageHandlePtrPtrGen), static_cast(&imageHandlePtrGen), - sizeof(sycl::ext::oneapi::experimental::unsampled_image_handle *)); + sizeof(syclexp::unsampled_image_handle *)); + q.wait_and_throw(); q.submit([&](sycl::handler &cgh) { - cgh.parallel_for(sycl::nd_range<1>{{width}, {width}}, [=](sycl::nd_item<1> - it) { - sycl::ext::oneapi::experimental::unsampled_image_handle * - *imageHandlePtrPtr = static_cast< - sycl::ext::oneapi::experimental::unsampled_image_handle **>( - imageHandlePtrPtrGen); - sycl::ext::oneapi::experimental::unsampled_image_handle *imageHandlePtr = - static_cast< - sycl::ext::oneapi::experimental::unsampled_image_handle *>( - imageHandlePtrPtr[0]); - sycl::ext::oneapi::experimental::unsampled_image_handle imageHandle = - imageHandlePtr[0]; - - size_t dim0 = it.get_local_id(0); - // Extension: read image data from handle - float pixel = sycl::ext::oneapi::experimental::fetch_image( - imageHandle, int(dim0)); - - // Extension: write to image data using handle - sycl::ext::oneapi::experimental::write_image(imgOut, int(dim0), pixel); - }); + cgh.parallel_for( + sycl::nd_range<1>{{width}, {width}}, [=](sycl::nd_item<1> it) { + syclexp::unsampled_image_handle **imageHandlePtrPtr = + static_cast( + imageHandlePtrPtrGen); + // Dereference the generic pointer to the unsampled_image_handle + // pointer + syclexp::unsampled_image_handle *imageHandlePtr = + static_cast( + imageHandlePtrPtr[0]); + // Dereference the unsampled_image_handle pointer + syclexp::unsampled_image_handle imageHandle = imageHandlePtr[0]; + + size_t dim0 = it.get_local_id(0); + // Extension: read image data from handle + float pixel = syclexp::fetch_image(imageHandle, int(dim0)); + + // Extension: write to image data using handle + syclexp::write_image(imgOut, int(dim0), pixel); + }); }); q.wait_and_throw(); @@ -82,8 +103,8 @@ int main() { q.wait_and_throw(); // Cleanup - sycl::ext::oneapi::experimental::destroy_image_handle(imgIn, q); - sycl::ext::oneapi::experimental::destroy_image_handle(imgOut, q); + syclexp::destroy_image_handle(imgIn, q); + syclexp::destroy_image_handle(imgOut, q); sycl::free(imageHandlePtrGen, q); sycl::free(imageHandlePtrPtrGen, q); From 6bb53731eda3900b30f10c6fbd8c69767f45888f Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Thu, 29 May 2025 15:25:54 +0100 Subject: [PATCH 4/7] Fix format Signed-off-by: JackAKirk --- .../test-e2e/bindless_images/array/fetch_handle_carray2d.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp index 633d9de503a01..714fbf8a5b01d 100644 --- a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp +++ b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp @@ -51,8 +51,9 @@ int main() { void *imageHandlePtrGen = sycl::malloc_device(sizeof(syclexp::unsampled_image_handle), q); - // Copy the create_image returned device unsampled_image_handle to the contents - // of the void* pointing to the manually created unsampled_image_handle + // Copy the create_image returned device unsampled_image_handle to the + // contents of the void* pointing to the manually created + // unsampled_image_handle q.memcpy(static_cast(imageHandlePtrGen), static_cast(&imgIn), sizeof(syclexp::unsampled_image_handle)); From ad7e3129be350899af56e5e5a833872b001de4dd Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Wed, 4 Jun 2025 17:13:32 +0100 Subject: [PATCH 5/7] Remove l0 XFAIL Signed-off-by: JackAKirk --- sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp index 714fbf8a5b01d..ffe9fe4ce8521 100644 --- a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp +++ b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp @@ -1,6 +1,4 @@ // REQUIRES: aspect-ext_oneapi_bindless_images -// XFAIL: level_zero -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18727 // UNSUPPORTED: hip // UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test. From 47bcc0a366abafecd9d48e162fe464ec7b6302b6 Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Thu, 5 Jun 2025 15:43:03 +0100 Subject: [PATCH 6/7] Xfail windows switch to sampled_image_handle example Signed-off-by: JackAKirk --- .../array/fetch_handle_carray2d.cpp | 45 ++++++++++--------- 1 file changed, 23 insertions(+), 22 deletions(-) diff --git a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp index ffe9fe4ce8521..6209465f13852 100644 --- a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp +++ b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp @@ -1,4 +1,6 @@ // REQUIRES: aspect-ext_oneapi_bindless_images +// XFAIL: level_zero && windows +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18727 // UNSUPPORTED: hip // UNSUPPORTED-INTENDED: Undetermined issue in 'create_image' in this test. @@ -34,55 +36,54 @@ int main() { q.ext_oneapi_copy(dataIn.data(), imgMemoryIn.get_handle(), desc); q.wait_and_throw(); // Extension: create the image and return the handle - syclexp::unsampled_image_handle imgIn = + syclexp::sampled_image_handle imgIn = syclexp::create_image(imgMemoryIn, desc, q); - syclexp::unsampled_image_handle imgOut = + syclexp::sampled_image_handle imgOut = syclexp::create_image(imgMemoryOut, desc, q); - // Copy the input data to the image_mem of the device unsampled_image_handle + // Copy the input data to the image_mem of the device sampled_image_handle q.ext_oneapi_copy(dataIn.data(), imgMemoryIn.get_handle(), desc); q.wait_and_throw(); - // Allocate an unsampled_image_handle manually instead of using create_image - // The implied purpose of this appears to be that a valid device void* is - // returned which points to the image_handle + // Allocate an sampled_image_handle manually instead of using create_image so + // we can allocate it on heap void *imageHandlePtrGen = - sycl::malloc_device(sizeof(syclexp::unsampled_image_handle), q); + sycl::malloc_device(sizeof(syclexp::sampled_image_handle), q); - // Copy the create_image returned device unsampled_image_handle to the - // contents of the void* pointing to the manually created - // unsampled_image_handle + // Copy the create_image returned device sampled_image_handle to the + // contents of the void* pointing to the heap allocated + // sampled_image_handle q.memcpy(static_cast(imageHandlePtrGen), static_cast(&imgIn), - sizeof(syclexp::unsampled_image_handle)); + sizeof(syclexp::sampled_image_handle)); q.wait_and_throw(); - // Allocate a device generic pointer pointing to an unsampled_image_handle* + // Allocate a device generic pointer pointing to an sampled_image_handle* void *imageHandlePtrPtrGen = - sycl::malloc_device(sizeof(syclexp::unsampled_image_handle *), q); + sycl::malloc_device(sizeof(syclexp::sampled_image_handle *), q); - // Copy the address of the manually allocated unsampled_image_handle to the + // Copy the address of the manually allocated sampled_image_handle to the // contents of the generic device pointer allocated above q.memcpy(static_cast(imageHandlePtrPtrGen), static_cast(&imageHandlePtrGen), - sizeof(syclexp::unsampled_image_handle *)); + sizeof(syclexp::sampled_image_handle *)); q.wait_and_throw(); q.submit([&](sycl::handler &cgh) { cgh.parallel_for( sycl::nd_range<1>{{width}, {width}}, [=](sycl::nd_item<1> it) { - syclexp::unsampled_image_handle **imageHandlePtrPtr = - static_cast( + syclexp::sampled_image_handle **imageHandlePtrPtr = + static_cast( imageHandlePtrPtrGen); - // Dereference the generic pointer to the unsampled_image_handle + // Dereference the generic pointer to the sampled_image_handle // pointer - syclexp::unsampled_image_handle *imageHandlePtr = - static_cast( + syclexp::sampled_image_handle *imageHandlePtr = + static_cast( imageHandlePtrPtr[0]); - // Dereference the unsampled_image_handle pointer - syclexp::unsampled_image_handle imageHandle = imageHandlePtr[0]; + // Dereference the sampled_image_handle pointer + syclexp::sampled_image_handle imageHandle = imageHandlePtr[0]; size_t dim0 = it.get_local_id(0); // Extension: read image data from handle From 2308810cd04419cc63e9d31c3b3663a7120e03bf Mon Sep 17 00:00:00 2001 From: JackAKirk Date: Wed, 11 Jun 2025 10:08:47 +0100 Subject: [PATCH 7/7] Use unsampled_image_handle Signed-off-by: JackAKirk --- .../array/fetch_handle_carray2d.cpp | 40 +++++++++---------- 1 file changed, 20 insertions(+), 20 deletions(-) diff --git a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp index 6209465f13852..0839d65c99a58 100644 --- a/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp +++ b/sycl/test-e2e/bindless_images/array/fetch_handle_carray2d.cpp @@ -36,54 +36,54 @@ int main() { q.ext_oneapi_copy(dataIn.data(), imgMemoryIn.get_handle(), desc); q.wait_and_throw(); // Extension: create the image and return the handle - syclexp::sampled_image_handle imgIn = + syclexp::unsampled_image_handle imgIn = syclexp::create_image(imgMemoryIn, desc, q); - syclexp::sampled_image_handle imgOut = + syclexp::unsampled_image_handle imgOut = syclexp::create_image(imgMemoryOut, desc, q); - // Copy the input data to the image_mem of the device sampled_image_handle + // Copy the input data to the image_mem of the device unsampled_image_handle q.ext_oneapi_copy(dataIn.data(), imgMemoryIn.get_handle(), desc); q.wait_and_throw(); - // Allocate an sampled_image_handle manually instead of using create_image so - // we can allocate it on heap + // Allocate an unsampled_image_handle manually instead of using create_image + // so we can allocate it on the heap void *imageHandlePtrGen = - sycl::malloc_device(sizeof(syclexp::sampled_image_handle), q); + sycl::malloc_device(sizeof(syclexp::unsampled_image_handle), q); - // Copy the create_image returned device sampled_image_handle to the + // Copy the create_image returned device unsampled_image_handle to the // contents of the void* pointing to the heap allocated - // sampled_image_handle + // unsampled_image_handle q.memcpy(static_cast(imageHandlePtrGen), static_cast(&imgIn), - sizeof(syclexp::sampled_image_handle)); + sizeof(syclexp::unsampled_image_handle)); q.wait_and_throw(); - // Allocate a device generic pointer pointing to an sampled_image_handle* + // Allocate a device generic pointer pointing to an unsampled_image_handle* void *imageHandlePtrPtrGen = - sycl::malloc_device(sizeof(syclexp::sampled_image_handle *), q); + sycl::malloc_device(sizeof(syclexp::unsampled_image_handle *), q); - // Copy the address of the manually allocated sampled_image_handle to the + // Copy the address of the manually allocated unsampled_image_handle to the // contents of the generic device pointer allocated above q.memcpy(static_cast(imageHandlePtrPtrGen), static_cast(&imageHandlePtrGen), - sizeof(syclexp::sampled_image_handle *)); + sizeof(syclexp::unsampled_image_handle *)); q.wait_and_throw(); q.submit([&](sycl::handler &cgh) { cgh.parallel_for( sycl::nd_range<1>{{width}, {width}}, [=](sycl::nd_item<1> it) { - syclexp::sampled_image_handle **imageHandlePtrPtr = - static_cast( + syclexp::unsampled_image_handle **imageHandlePtrPtr = + static_cast( imageHandlePtrPtrGen); - // Dereference the generic pointer to the sampled_image_handle + // Dereference the generic pointer to the unsampled_image_handle // pointer - syclexp::sampled_image_handle *imageHandlePtr = - static_cast( + syclexp::unsampled_image_handle *imageHandlePtr = + static_cast( imageHandlePtrPtr[0]); - // Dereference the sampled_image_handle pointer - syclexp::sampled_image_handle imageHandle = imageHandlePtr[0]; + // Dereference the unsampled_image_handle pointer + syclexp::unsampled_image_handle imageHandle = imageHandlePtr[0]; size_t dim0 = it.get_local_id(0); // Extension: read image data from handle