[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

@@ -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 b766009add8dc41ad03e2a63be9f6ed40eea4c55 (HEAD, origin/main, origin/HEAD)
- # Merge: 608e9184 c3c57284
+ set(UNIFIED_RUNTIME_REPO "https://github.com/Seanst98/unified-runtime.git")
+ # commit d0a50523006fa6f283da6a36811081add3bb22fc
+ # Merge: 804851e4 04deb8b3
  # Author: Omar Ahmed <omar.ahmed@codeplay.com>
- # Date: Tue Aug 27 16:30:04 2024 +0100
- # Merge pull request #2011 from nrspruit/fix_opencl_usm_align_check
- # [OpenCL] Fix USM alignment error check to occur always and return nulllptr
- set(UNIFIED_RUNTIME_TAG b766009add8dc41ad03e2a63be9f6ed40eea4c55)
+ # 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 54fcde0eb5874a558682d60e248d680182200847)
 
  set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
  # Due to the use of dependentloadflag and no installer for UMF and hwloc we need

sycl/test-e2e/bindless_images/read_norm_types.cpp

@@ -6,6 +6,7 @@
 #include <iostream>
 #include <limits>
 #include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
 
 #include "helpers/common.hpp"
 #include <sycl/ext/oneapi/bindless_images.hpp>
@@ -29,7 +30,7 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
 
  std::vector<InputType> dataIn(numElems, InputType((DType)dtypeMaxVal));
  std::vector<OutputType> dataOut(numElems);
- std::vector<OutputType> expected(numElems, OutputType(1.f));
+ std::vector<OutputType> expected(numElems, OutputType(2.f));
 
  try {
 
@@ -47,9 +48,30 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> 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,27 +82,35 @@ bool run_test(sycl::range<NDims> globalSize, sycl::range<NDims> localSize) {
  if constexpr (NDims == 1) {
  size_t dim0 = it.get_global_id(0);
  float fdim0 = dim0 / globalSize[0];
- OutputType pixel =
- syclexp::sample_image<OutputType>(imgIn, fdim0);
- syclexp::write_image(imgOut, int(dim0), pixel);
+ OutputType pixel1 =
+ syclexp::sample_image<OutputType>(imgIn1, fdim0);
+ OutputType pixel2 =
+ syclexp::sample_image<OutputType>(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<OutputType>(
- imgIn, sycl::float2(fdim0, fdim1));
- syclexp::write_image(imgOut, sycl::int2(dim0, dim1), pixel);
+ OutputType pixel1 = syclexp::sample_image<OutputType>(
+ imgIn1, sycl::float2(fdim0, fdim1));
+ OutputType pixel2 = syclexp::sample_image<OutputType>(
+ 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);
  size_t dim2 = it.get_global_id(2);
  float fdim0 = dim0 / globalSize[0];
  float fdim1 = dim1 / globalSize[1];
  float fdim2 = dim2 / globalSize[2];
- OutputType pixel = syclexp::sample_image<OutputType>(
- imgIn, sycl::float3(fdim0, fdim1, fdim2));
- syclexp::write_image(imgOut, sycl::int3(dim0, dim1, dim2), pixel);
+ OutputType pixel1 = syclexp::sample_image<OutputType>(
+ imgIn1, sycl::float3(fdim0, fdim1, fdim2));
+ OutputType pixel2 = syclexp::sample_image<OutputType>(
+ 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<NDims> globalSize, sycl::range<NDims> 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;