From af9429698ae03f551b7c0c4544ee7b9babbb7afb Mon Sep 17 00:00:00 2001 From: Sean Stirling Date: Fri, 23 Aug 2024 15:10:33 +0100 Subject: [PATCH] [SYCL][Bindless][E2E] Test normalized usm bindless images Modify pre-existing read_norm_types test to add testing for 2D USM bindless images. --- sycl/cmake/modules/FetchUnifiedRuntime.cmake | 16 ++--- .../bindless_images/read_norm_types.cpp | 62 +++++++++++++++---- 2 files changed, 58 insertions(+), 20 deletions(-) diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 956c33bec68df..cc08305676b30 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -116,14 +116,14 @@ if(SYCL_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") - # commit 8c9dd7e464a99ebbfb238ac2dabefc3ac77baea5 - # Merge: a99dbcee 3abe18cf - # Author: Piotr Balcer - # Date: Fri Sep 6 17:21:17 2024 +0200 - # Merge pull request #1820 from pbalcer/static-linking - # Add support for static linking of the L0 adapter - set(UNIFIED_RUNTIME_TAG 8c9dd7e464a99ebbfb238ac2dabefc3ac77baea5) + set(UNIFIED_RUNTIME_REPO "https://github.com/Seanst98/unified-runtime.git") + # commit d0a50523006fa6f283da6a36811081add3bb22fc + # Merge: 804851e4 04deb8b3 + # Author: Omar Ahmed + # Date: Tue Aug 20 16:28:30 2024 +0100 + # Merge pull request #1940 from RossBrunton/ross/urcall + # [XPTI] Use `ur.call` rather than `ur` in XPTI + set(UNIFIED_RUNTIME_TAG 94cb7b07e5dc5712432d9793b2879916dc9b8653) 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;