[SYCL][Bindless][E2E] Test normalized 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 22962057df1b9d538e08088a7b75d9d8e7c29f90 (HEAD, origin/main, origin/HEAD)
-  # Merge: e824ddc2 f0a1c433
-  # Author: aarongreig <aaron.greig@codeplay.com>
-  # Date:   Fri Sep 27 16:54:04 2024 +0100
-  #     Merge pull request #2017 from nrspruit/new_sysman_init
-  #     [L0] Use zesInit for SysMan API usage
-  set(UNIFIED_RUNTIME_TAG 22962057df1b9d538e08088a7b75d9d8e7c29f90)
+  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 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 bcf2244dccdef352afaf4d4520526573876981e3)
 
   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;