diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index ee302b87e4b01..bec7c8ffba0da 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -117,14 +117,14 @@ if(SYCL_UR_USE_FETCH_CONTENT) endfunction() set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") - # commit 00f958f375205fd86309f95b925141cf664ff955 - # Merge: cc2d5909 98a67a2e + # commit 1f13d2ceb0494d84ce7b32f6b453dbb256fb702a + # Merge: 5276c534 bcf2244d # Author: aarongreig - # Date: Wed Oct 2 09:51:21 2024 +0100 - # Merge pull request #2139 from nrspruit/zeHandle_copy_dependencies - # [L0] Pass and track event dependencies required before executing Memory - # Copy buffer inits - set(UNIFIED_RUNTIME_TAG 00f958f375205fd86309f95b925141cf664ff955) + # Date: Wed Oct 2 15:04:33 2024 +0100 + # Merge pull request #2056 from Seanst98/sean/usm-normalized-fix + # + # [CUDA][Bindless] Address USM normalized type image creation failure and functionality + set(UNIFIED_RUNTIME_TAG 1f13d2ceb0494d84ce7b32f6b453dbb256fb702a) set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES") # Due to the use of dependentloadflag and no installer for UMF and hwloc we need diff --git a/sycl/test-e2e/bindless_images/read_norm_types.cpp b/sycl/test-e2e/bindless_images/read_norm_types.cpp index 16de22f2f69ce..ddc185b71d720 100644 --- a/sycl/test-e2e/bindless_images/read_norm_types.cpp +++ b/sycl/test-e2e/bindless_images/read_norm_types.cpp @@ -6,6 +6,7 @@ #include #include #include +#include #include "helpers/common.hpp" #include @@ -29,7 +30,7 @@ bool run_test(sycl::range globalSize, sycl::range localSize) { std::vector dataIn(numElems, InputType((DType)dtypeMaxVal)); std::vector dataOut(numElems); - std::vector expected(numElems, OutputType(1.f)); + std::vector expected(numElems, OutputType(2.f)); try { @@ -47,9 +48,30 @@ bool run_test(sycl::range globalSize, sycl::range localSize) { sycl::coordinate_normalization_mode::normalized, sycl::filtering_mode::nearest}; - auto imgIn = syclexp::create_image(imgMemIn, sampler, descIn, q); + auto imgIn1 = syclexp::create_image(imgMemIn, sampler, descIn, q); auto imgOut = syclexp::create_image(imgMemOut, descOut, q); + void *allocUSM = nullptr; + syclexp::image_mem_handle allocMem; + syclexp::sampled_image_handle imgIn2; + + if constexpr (NDims == 2) { + size_t pitch = 0; + allocUSM = syclexp::pitched_alloc_device(&pitch, descIn, q); + + if (allocUSM == nullptr) { + std::cerr << "Error allocating 2D USM memory!" << std::endl; + return false; + } + imgIn2 = syclexp::create_image(allocUSM, pitch, sampler, descIn, q); + q.ext_oneapi_copy(dataIn.data(), allocUSM, descIn, pitch); + + } else { + allocMem = syclexp::alloc_image_mem(descIn, q); + imgIn2 = syclexp::create_image(allocMem, sampler, descIn, q); + q.ext_oneapi_copy(dataIn.data(), allocMem, descIn); + } + q.ext_oneapi_copy(dataIn.data(), imgMemIn, descIn); q.wait_and_throw(); @@ -60,17 +82,22 @@ bool run_test(sycl::range globalSize, sycl::range localSize) { if constexpr (NDims == 1) { size_t dim0 = it.get_global_id(0); float fdim0 = dim0 / globalSize[0]; - OutputType pixel = - syclexp::sample_image(imgIn, fdim0); - syclexp::write_image(imgOut, int(dim0), pixel); + OutputType pixel1 = + syclexp::sample_image(imgIn1, fdim0); + OutputType pixel2 = + syclexp::sample_image(imgIn2, fdim0); + syclexp::write_image(imgOut, int(dim0), pixel1 + pixel2); } else if constexpr (NDims == 2) { size_t dim0 = it.get_global_id(0); size_t dim1 = it.get_global_id(1); float fdim0 = dim0 / globalSize[0]; float fdim1 = dim1 / globalSize[1]; - OutputType pixel = syclexp::sample_image( - imgIn, sycl::float2(fdim0, fdim1)); - syclexp::write_image(imgOut, sycl::int2(dim0, dim1), pixel); + OutputType pixel1 = syclexp::sample_image( + imgIn1, sycl::float2(fdim0, fdim1)); + OutputType pixel2 = syclexp::sample_image( + imgIn2, sycl::float2(fdim0, fdim1)); + syclexp::write_image(imgOut, sycl::int2(dim0, dim1), + pixel1 + pixel2); } else if constexpr (NDims == 3) { size_t dim0 = it.get_global_id(0); size_t dim1 = it.get_global_id(1); @@ -78,9 +105,12 @@ bool run_test(sycl::range globalSize, sycl::range localSize) { float fdim0 = dim0 / globalSize[0]; float fdim1 = dim1 / globalSize[1]; float fdim2 = dim2 / globalSize[2]; - OutputType pixel = syclexp::sample_image( - imgIn, sycl::float3(fdim0, fdim1, fdim2)); - syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2), pixel); + OutputType pixel1 = syclexp::sample_image( + imgIn1, sycl::float3(fdim0, fdim1, fdim2)); + OutputType pixel2 = syclexp::sample_image( + imgIn2, sycl::float3(fdim0, fdim1, fdim2)); + syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2), + pixel1 + pixel2); } }); }); @@ -89,12 +119,20 @@ bool run_test(sycl::range globalSize, sycl::range localSize) { q.ext_oneapi_copy(imgMemOut, dataOut.data(), descOut); q.wait_and_throw(); - syclexp::destroy_image_handle(imgIn, q); + syclexp::destroy_image_handle(imgIn1, q); + syclexp::destroy_image_handle(imgIn2, q); syclexp::destroy_image_handle(imgOut, q); syclexp::free_image_mem(imgMemIn, syclexp::image_type::standard, dev, ctxt); syclexp::free_image_mem(imgMemOut, syclexp::image_type::standard, dev, ctxt); + + if constexpr (NDims == 2) { + sycl::free(allocUSM, ctxt); + } else { + syclexp::free_image_mem(allocMem, syclexp::image_type::standard, dev, + ctxt); + } } catch (sycl::exception e) { std::cerr << "SYCL exception caught! : " << e.what() << "\n"; return false;