[sycl-rel] Cherry-pick sycl patches

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -89,6 +89,19 @@ static const char *LegalSYCLFunctionsInStatelessMode[] = {
 
 namespace {
 
+class BuffDeleter {
+public:
+  BuffDeleter(char *Buffer) : Buff(Buffer) {};
+  ~BuffDeleter() { std::free(Buff); };
+
+  BuffDeleter() = delete;
+  BuffDeleter(const BuffDeleter &) = delete;
+  BuffDeleter(BuffDeleter &&) = delete;
+
+private:
+  char *Buff;
+};
+
 class ESIMDVerifierImpl {
   const Module &M;
   bool MayNeedForceStatelessMemModeAPI;
@@ -150,6 +163,7 @@ class ESIMDVerifierImpl {
 
           id::OutputBuffer NameBuf;
           NameNode->print(NameBuf);
+          BuffDeleter NameBufDeleter(NameBuf.getBuffer());
           StringRef Name(NameBuf.getBuffer(), NameBuf.getCurrentPosition());
 
           // We are interested in functions defined in SYCL namespace, but

sycl/source/detail/device_impl.cpp

@@ -671,17 +671,15 @@ bool device_impl::has(aspect Aspect) const {
     ur_bool_t support = false;
     bool call_successful =
         getAdapter()->call_nocheck<UrApiKind::urDeviceGetInfo>(
-            MDevice,
-            UR_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_1D_USM_SUPPORT_EXP,
+            MDevice, UR_DEVICE_INFO_BINDLESS_SAMPLE_1D_USM_SUPPORT_EXP,
             sizeof(ur_bool_t), &support, nullptr) == UR_RESULT_SUCCESS;
     return call_successful && support;
   }
   case aspect::ext_oneapi_bindless_images_sample_2d_usm: {
     ur_bool_t support = false;
     bool call_successful =
         getAdapter()->call_nocheck<UrApiKind::urDeviceGetInfo>(
-            MDevice,
-            UR_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_2D_USM_SUPPORT_EXP,
+            MDevice, UR_DEVICE_INFO_BINDLESS_SAMPLE_2D_USM_SUPPORT_EXP,
             sizeof(ur_bool_t), &support, nullptr) == UR_RESULT_SUCCESS;
     return call_successful && support;
   }

sycl/test-e2e/DeviceLib/assert-windows.cpp

@@ -1,22 +1,15 @@
-// REQUIRES: cpu,windows
+// REQUIRES: windows
+// XFAIL: opencl && gpu
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/11364
 //
-// RUN: %{build} -o %t.out
+// RUN: %{build} -DSYCL_FALLBACK_ASSERT=1 -o %t.out
 //
-// MSVC implementation of assert does not call an unreachable built-in, so the
-// program doesn't terminate when fallback is used.
-//
-// FIXME: SPIR-V Unreachable should be called from the fallback
-// explicitly. Since the test is going to crash, we'll have to follow a similar
-// approach as on Linux - call the test in a subprocess.
-//
-// RUN: env SYCL_UR_TRACE=2 SYCL_DEVICELIB_INHIBIT_NATIVE=1 CL_CONFIG_USE_VECTORIZER=False %{run} %t.out | FileCheck %s --check-prefix=CHECK-FALLBACK
-// RUN: env SHOULD_CRASH=1 SYCL_DEVICELIB_INHIBIT_NATIVE=1 CL_CONFIG_USE_VECTORIZER=False %{run} %t.out | FileCheck %s --check-prefix=CHECK-MESSAGE
+// RUN: not env SHOULD_CRASH=1 SYCL_DEVICELIB_INHIBIT_NATIVE=1 CL_CONFIG_USE_VECTORIZER=False \
+// RUN: %{run} %t.out 2>&1 >/dev/null | FileCheck %s --check-prefix=CHECK-MESSAGE
 //
 // CHECK-MESSAGE: {{.*}}assert-windows.cpp:{{[0-9]+}}: (null): global id:
 // [{{[0-3]}},0,0], local id: [{{[0-3]}},0,0] Assertion `accessorC[wiID] == 0 &&
 // "Invalid value"` failed.
-//
-// CHECK-FALLBACK: <--- urProgramLink
 
 #include "../helpers.hpp"
 #include <array>

sycl/test-e2e/Matrix/SG32/get_coordinate_ops.cpp

@@ -15,9 +15,6 @@
 // REQUIRES: aspect-ext_intel_matrix
 // REQUIRES-INTEL-DRIVER: lin: 30049, win: 101.4943
 
-// XFAIL: arch-intel_gpu_pvc
-// XFAIL-TRACKER: GSD-10524
-
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/Matrix/SG32/joint_matrix_annotated_ptr.cpp

@@ -9,12 +9,13 @@
 // UNSUPPORTED-INTENDED: aspect-ext_intel_matrix isn't currently supported for
 // other triples
 
+// UNSUPPORTED: gpu-intel-dg2
+// UNSUPPORTED-INTENDED: SG size = 32 is not currently supported for SYCL Joint
+// Matrix by IGC on DG2
+
 // REQUIRES: aspect-ext_intel_matrix
 // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943
 
-// XFAIL: run-mode && gpu-intel-dg2
-// XFAIL-TRACKER: GSD-4181
-
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/Matrix/SG32/joint_matrix_colA_rowB_colC.cpp

@@ -9,12 +9,13 @@
 // UNSUPPORTED-INTENDED: aspect-ext_intel_matrix isn't currently supported for
 // other triples
 
+// UNSUPPORTED: gpu-intel-dg2
+// UNSUPPORTED-INTENDED: SG size = 32 is not currently supported for SYCL Joint
+// Matrix by IGC on DG2
+
 // REQUIRES: aspect-ext_intel_matrix
 // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943
 
-// XFAIL: run-mode && gpu-intel-dg2
-// XFAIL-TRACKER: GSD-5768
-
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/Matrix/SG32/joint_matrix_int8_colmajorA_colmajorB.cpp

@@ -16,11 +16,14 @@
 // RUN: %{run} %t.out
 
 // This tests support of col major layout for matrix B which does transpose and
-// then VNNI transform. This is currently only available on AMX
+// then VNNI transform. This is currently only available on AMX and PVC
 
-// XFAIL: gpu
+// XFAIL: arch-intel_gpu_bmg_g21
 // XFAIL-TRACKER: GSD-5768
 
+// UNSUPPORTED: gpu-intel-dg2
+// UNSUPPORTED-INTENDED: SG size = 32 is not currently supported for SYCL Joint Matrix by IGC on DG2
+
 #include "common.hpp"
 
 using namespace sycl;

sycl/test-e2e/Matrix/joint_matrix_int8_colmajorA_colmajorB.cpp

@@ -15,9 +15,9 @@
 // RUN: %{run} %t.out
 
 // This tests support of col major layout for matrix B which does transpose and
-// then VNNI transform. This is currently only available on AMX
+// then VNNI transform. This is currently only available on AMX and PVC
 
-// XFAIL: gpu
+// XFAIL: gpu-intel-dg2 || arch-intel_gpu_bmg_g21
 // XFAIL-TRACKER: GSD-5768
 
 #include "common.hpp"

sycl/test-e2e/bindless_images/sampled_fetch/fetch_1D_USM_host.cpp

@@ -0,0 +1,102 @@
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_bindless_sampled_image_fetch_1d_usm
+// UNSUPPORTED: target-amd
+// UNSUPPORTED-INTENDED: Sampled fetch not currently supported on AMD
+
+// RUN: %{build} -o %t.out
+// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/bindless_images.hpp>
+#include <sycl/usm.hpp>
+
+class kernel_sampled_fetch;
+
+// Uncomment to print additional test information
+// #define VERBOSE_PRINT
+
+int main() {
+  sycl::device dev;
+  sycl::queue q(dev);
+  auto ctxt = q.get_context();
+
+  // Declare image size, and expected output and actual output vectors
+  constexpr size_t width = 32;
+  constexpr size_t widthInBytes = width * sizeof(float);
+  std::vector<float> out(width);
+  std::vector<float> expected(width);
+  for (int i = 0; i < width; ++i) {
+    expected[i] = static_cast<float>(i);
+  }
+
+  namespace syclexp = sycl::ext::oneapi::experimental;
+
+  try {
+    // Extension: image descriptor
+    syclexp::image_descriptor desc({width}, 1, sycl::image_channel_type::fp32);
+
+    // Extension: Image creation requires a sampler, but it will have no effect
+    //            on the result, as we will use `fetch_image` in the kernel.
+    syclexp::bindless_image_sampler samp(
+        sycl::addressing_mode::repeat,
+        sycl::coordinate_normalization_mode::normalized,
+        sycl::filtering_mode::linear);
+
+    // Allocate Host USM and initialize with expected data
+    float *imgMem = sycl::malloc_host<float>(width, q);
+    memcpy(imgMem, expected.data(), widthInBytes);
+
+    // Extension: create the image backed by Host USM and return the handle
+    auto imgHandle = syclexp::create_image(imgMem, 0, samp, desc, q);
+
+    // Create a buffer to output the result from `fetch_image`
+    sycl::buffer outBuf(out.data(), sycl::range{width});
+    q.submit([&](sycl::handler &cgh) {
+      sycl::accessor outAcc{outBuf, cgh, sycl::write_only};
+
+      cgh.parallel_for<kernel_sampled_fetch>(width, [=](sycl::id<1> id) {
+        // Extension: fetch data from sampled image handle
+        outAcc[id] = syclexp::fetch_image<float>(imgHandle, int(id[0]));
+      });
+    });
+
+    q.wait_and_throw();
+
+    // Extension: cleanup
+    syclexp::destroy_image_handle(imgHandle, dev, ctxt);
+    sycl::free(imgMem, ctxt);
+  } catch (sycl::exception e) {
+    std::cerr << "SYCL exception caught! : " << e.what() << "\n";
+    return 1;
+  } catch (...) {
+    std::cerr << "Unknown exception caught!\n";
+    return 2;
+  }
+
+  // collect and validate output
+  bool validated = true;
+  for (int i = 0; i < width; i++) {
+    bool mismatch = false;
+    if (out[i] != expected[i]) {
+      mismatch = true;
+      validated = false;
+    }
+
+    if (mismatch) {
+#ifdef VERBOSE_PRINT
+      std::cout << "Result mismatch! Expected: " << expected[i]
+                << ", Actual: " << out[i] << std::endl;
+#else
+      break;
+#endif
+    }
+  }
+  if (validated) {
+    std::cout << "Test passed!" << std::endl;
+    return 0;
+  }
+
+  std::cout << "Test failed!" << std::endl;
+  return 3;
+}

sycl/test-e2e/bindless_images/sampling_1D_USM_host.cpp

@@ -0,0 +1,126 @@
+// REQUIRES: aspect-ext_oneapi_bindless_images
+// REQUIRES: aspect-ext_oneapi_bindless_images_sample_1d_usm
+
+// UNSUPPORTED: hip
+// UNSUPPORTED-INTENDED: Host USM backed image support is not yet enabled in UR
+// adapter. Also, when provionally enabled, the test crashes upon image
+// creation, whereas Device USM backed images do not crash. This issue is
+// undetermined.
+
+// RUN: %{build} -o %t.out
+// RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
+
+#include <cmath>
+#include <iostream>
+#include <sycl/detail/core.hpp>
+
+#include <sycl/ext/oneapi/bindless_images.hpp>
+#include <sycl/usm.hpp>
+
+// Uncomment to print additional test information
+// #define VERBOSE_PRINT
+
+class sample_host_usm_image_kernel;
+
+int main() {
+
+  sycl::device dev;
+  sycl::queue q(dev);
+  auto ctxt = q.get_context();
+
+  // declare image data
+  size_t width = 32;
+  size_t widthInBytes = width * sizeof(float);
+  std::vector<float> out(width);
+  std::vector<float> expected(width);
+  for (int i = 0; i < width; ++i) {
+    expected[i] = static_cast<float>(i);
+  }
+
+  try {
+    sycl::ext::oneapi::experimental::bindless_image_sampler samp(
+        sycl::addressing_mode::clamp,
+        sycl::coordinate_normalization_mode::normalized,
+        sycl::filtering_mode::linear);
+
+    // Extension: image descriptor
+    sycl::ext::oneapi::experimental::image_descriptor desc(
+        {width}, 1, sycl::image_channel_type::fp32);
+
+    // Host USM allocation
+    float *imgMem = sycl::malloc_host<float>(width, ctxt);
+
+    if (imgMem == nullptr) {
+      std::cerr << "Error allocating host USM!" << std::endl;
+      return 1;
+    }
+
+    // Initialize input data
+    for (int i = 0; i < width; ++i) {
+      imgMem[i] = static_cast<float>(i);
+    }
+
+    // Extension: create the image and return the handle
+    sycl::ext::oneapi::experimental::sampled_image_handle imgHandle =
+        sycl::ext::oneapi::experimental::create_image(imgMem, 0 /* pitch */,
+                                                      samp, desc, dev, ctxt);
+
+    sycl::buffer<float, 1> buf((float *)out.data(), sycl::range<1>{width});
+    q.submit([&](sycl::handler &cgh) {
+      auto outAcc =
+          buf.get_access<sycl::access_mode::write>(cgh, sycl::range<1>{width});
+
+      cgh.parallel_for<sample_host_usm_image_kernel>(
+          sycl::nd_range<1>{{width}, {width}}, [=](sycl::nd_item<1> it) {
+            size_t dim0 = it.get_local_id(0);
+
+            // Normalize coordinates -- +0.5 to look towards centre of pixel
+            float fdim0 = float(dim0 + 0.5f) / (float)width;
+
+            // Extension: sample image data from handle
+            float px = sycl::ext::oneapi::experimental::sample_image<float>(
+                imgHandle, (float)fdim0);
+
+            outAcc[sycl::id<1>{dim0}] = px;
+          });
+    });
+
+    q.wait_and_throw();
+
+    // Extension: cleanup
+    sycl::ext::oneapi::experimental::destroy_image_handle(imgHandle, dev, ctxt);
+    sycl::free(imgMem, ctxt);
+  } catch (sycl::exception e) {
+    std::cerr << "SYCL exception caught! : " << e.what() << "\n";
+    return 1;
+  } catch (...) {
+    std::cerr << "Unknown exception caught!\n";
+    return 2;
+  }
+
+  // collect and validate output
+  bool validated = true;
+  for (int i = 0; i < width; i++) {
+    bool mismatch = false;
+    if (out[i] != expected[i]) {
+      mismatch = true;
+      validated = false;
+    }
+
+    if (mismatch) {
+#ifdef VERBOSE_PRINT
+      std::cout << "Result mismatch! Expected: " << expected[i]
+                << ", Actual: " << out[i] << std::endl;
+#else
+      break;
+#endif
+    }
+  }
+  if (validated) {
+    std::cout << "Test passed!" << std::endl;
+    return 0;
+  }
+
+  std::cout << "Test failed!" << std::endl;
+  return 3;
+}

sycl/unittests/Extensions/CMakeLists.txt

@@ -29,3 +29,4 @@ add_subdirectory(VirtualFunctions)
 add_subdirectory(VirtualMemory)
 add_subdirectory(NumComputeUnits)
 add_subdirectory(FreeFunctionCommands)
+add_subdirectory(KernelQueries)

sycl/unittests/Extensions/KernelQueries/CMakeLists.txt

@@ -0,0 +1,3 @@
+add_sycl_unittest(KernelQueriesTests OBJECT
+  SpillMemorySize.cpp
+)