Merge remote-tracking branch 'upstream/sycl' into georgi/hip-undef-cuda-arch

Author: Georgi Mirazchiyski

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 <[email protected]>
-  # 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

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;