diff --git a/sycl/cmake/modules/UnifiedRuntimeTag.cmake b/sycl/cmake/modules/UnifiedRuntimeTag.cmake index 5ba8a6aa938f1..7b46bd5b034c9 100644 --- a/sycl/cmake/modules/UnifiedRuntimeTag.cmake +++ b/sycl/cmake/modules/UnifiedRuntimeTag.cmake @@ -1,7 +1,7 @@ -# commit b841691699393dd2375e987c3d38d5f59c3e35cf -# Merge: c6859445 9de10cd9 +# commit 0bb6789f0113ea937d861fd67fd677b91ecdeb8b +# Merge: e370a2b9 eeff9f4a # Author: Kenneth Benzie (Benie) -# Date: Thu Jan 23 16:07:06 2025 +0000 -# Merge pull request #2559 from Bensuo/fix_kernel_arg_indices -# [CUDA][HIP] Fix kernel arguments being overwritten when added out of order -set(UNIFIED_RUNTIME_TAG b841691699393dd2375e987c3d38d5f59c3e35cf) +# Date: Mon Jan 27 10:40:02 2025 +0000 +# Merge pull request #2551 from przemektmalon/przemek/bindless-images-host-usm +# Enable creation of bindless images backed by host USM +set(UNIFIED_RUNTIME_TAG 0bb6789f0113ea937d861fd67fd677b91ecdeb8b) diff --git a/sycl/test-e2e/bindless_images/sampling_2D_USM_host.cpp b/sycl/test-e2e/bindless_images/sampling_2D_USM_host.cpp new file mode 100644 index 0000000000000..bca3d2c1c0ddd --- /dev/null +++ b/sycl/test-e2e/bindless_images/sampling_2D_USM_host.cpp @@ -0,0 +1,147 @@ +// REQUIRES: cuda +// REQUIRES: aspect-ext_oneapi_bindless_images_2d_usm + +// RUN: %{build} -o %t.out +// RUN: %{run-unfiltered-devices} %t.out + +#include +#include +#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 + size_t width = 5; + size_t height = 6; + size_t N = width * height; + size_t widthInBytes = width * sizeof(float); + 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++) { + expected[i + (width * j)] = i + (width * j); + dataIn[i + (width * j)] = i + (width * j); + } + } + + try { + sycl::ext::oneapi::experimental::bindless_image_sampler samp( + sycl::addressing_mode::clamp, + sycl::coordinate_normalization_mode::normalized, + sycl::filtering_mode::linear); + + // Extension: image descriptor + sycl::ext::oneapi::experimental::image_descriptor desc( + {width, height}, 1, sycl::image_channel_type::fp32); + + auto devicePitchAlign = dev.get_info< + sycl::ext::oneapi::experimental::info::device::image_row_pitch_align>(); + auto deviceMaxPitch = + dev.get_info(); + + // Pitch requirements: + // - pitch % devicePitchAlign == 0 + // - pitch >= widthInBytes + // - pitch <= deviceMaxPitch + size_t pitch = devicePitchAlign * + std::ceil(float(widthInBytes) / float(devicePitchAlign)); + assert(pitch <= deviceMaxPitch); + + // Host USM allocation + auto imgMem = + sycl::aligned_alloc_host(devicePitchAlign, (pitch * height), ctxt); + + if (imgMem == nullptr) { + std::cerr << "Error allocating images!" << std::endl; + return 1; + } + + // Copy to host USM and incorporate pitch + for (size_t i = 0; i < height; i++) { + memcpy(static_cast(imgMem) + (i * pitch / sizeof(float)), + dataIn.data() + (i * width), widthInBytes); + } + + // Extension: create the image and return the handle + sycl::ext::oneapi::experimental::sampled_image_handle imgHandle = + sycl::ext::oneapi::experimental::create_image(imgMem, pitch, samp, desc, + dev, ctxt); + + sycl::buffer buf((float *)out.data(), + sycl::range<2>{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); + + // Normalize coordinates -- +0.5 to look towards centre of pixel + float fdim0 = float(dim0 + 0.5f) / (float)width; + float fdim1 = float(dim1 + 0.5f) / (float)height; + + // Extension: sample image data from handle + float px = sycl::ext::oneapi::experimental::sample_image( + imgHandle, sycl::float2(fdim0, fdim1)); + + outAcc[sycl::id<2>{dim1, dim0}] = px; + }); + }); + + q.wait_and_throw(); + + // Extension: cleanup + sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle, dev, ctxt); + sycl::free(imgMem, 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; +}