Merge from 'sycl' to 'sycl-web' (8 commits)

Author: iclsrc

.github/CODEOWNERS

@@ -139,18 +139,18 @@ devops/dependencies-igc-dev.json @intel/sycl-matrix-reviewers @intel/dpcpp-esimd
 # Benchmarking scripts
 devops/scripts/benchmarks/ @intel/llvm-reviewers-benchmarking
 
-# Kernel fusion JIT compiler
-sycl-jit/ @intel/dpcpp-kernel-fusion-reviewers
-sycl/doc/design/KernelFusionJIT.md @intel/dpcpp-kernel-fusion-reviewers
-sycl/doc/extensions/experimental/sycl_ext_codeplay_kernel_fusion.asciidoc @intel/dpcpp-kernel-fusion-reviewers
-sycl/include/sycl/ext/codeplay/experimental/fusion_properties.hpp  @intel/dpcpp-kernel-fusion-reviewers
-sycl/include/sycl/ext/codeplay/experimental/fusion_wrapper.hpp @intel/dpcpp-kernel-fusion-reviewers
-sycl/source/detail/fusion/ @intel/dpcpp-kernel-fusion-reviewers
-sycl/source/detail/jit_compiler.hpp @intel/dpcpp-kernel-fusion-reviewers
-sycl/source/detail/jit_compiler.cpp @intel/dpcpp-kernel-fusion-reviewers
-sycl/source/detail/jit_device_binaries.hpp @intel/dpcpp-kernel-fusion-reviewers
-sycl/source/detail/jit_device_binaries.cpp @intel/dpcpp-kernel-fusion-reviewers
-sycl/test-e2e/KernelFusion @intel/dpcpp-kernel-fusion-reviewers
+# JIT compiler
+sycl-jit/ @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
+sycl/doc/design/KernelFusionJIT.md @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
+sycl/doc/extensions/experimental/sycl_ext_codeplay_kernel_fusion.asciidoc @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
+sycl/include/sycl/ext/codeplay/experimental/fusion_properties.hpp  @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
+sycl/include/sycl/ext/codeplay/experimental/fusion_wrapper.hpp @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
+sycl/source/detail/fusion/ @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
+sycl/source/detail/jit_compiler.hpp @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
+sycl/source/detail/jit_compiler.cpp @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
+sycl/source/detail/jit_device_binaries.hpp @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
+sycl/source/detail/jit_device_binaries.cpp @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
+sycl/test-e2e/KernelFusion @intel/llvm-reviewers-runtime @cperkinsintel @aelovikov-intel
 
 # Matrix
 sycl/include/sycl/ext/oneapi/matrix/ @intel/sycl-matrix-reviewers

devops/dependencies-igc-dev.json

@@ -1,10 +1,10 @@
 {
   "linux": {
     "igc_dev": {
-      "github_tag": "igc-dev-4e160b3",
-      "version": "4e160b3",
-      "updated_at": "2025-09-11T02:09:40Z",
-      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/3981554499/zip",
+      "github_tag": "igc-dev-5a47189",
+      "version": "5a47189",
+      "updated_at": "2025-09-18T02:11:35Z",
+      "url": "https://api.github.com/repos/intel/intel-graphics-compiler/actions/artifacts/4041051864/zip",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     }
   }

sycl-jit/jit-compiler/include/Resource.h

@@ -8,13 +8,29 @@
 
 #pragma once
 
-#include <iterator>
-#include <string_view>
-#include <utility>
-
-namespace jit_compiler {
+namespace jit_compiler::resource {
+// `resource.cpp` is compiled using freshly built clang and it's very hard to
+// sync compilation options between that and normal compilation for other files.
+// Note that some of the options might affect ABI (e.g., libstdc++ vs. libc++
+// usage, or custom sysroot/gcc installation directory). A much easier approach
+// is to ensure that `resource.cpp` doesn't have any includes at all, hence
+// these helpers:
+template <class T, unsigned long long N>
+constexpr unsigned long long size(const T (&)[N]) noexcept {
+  return N;
+}
+struct resource_string_view {
+  template <unsigned long long N>
+  resource_string_view(const char (&S)[N]) : S(S), Size(N - 1) {}
+  const char *S;
+  unsigned long long Size;
+};
+struct resource_file {
+  resource_string_view Path;
+  resource_string_view Content;
+};
 // Defined in the auto-generated file:
-extern const std::pair<std::string_view, std::string_view> ToolchainFiles[];
-extern size_t NumToolchainFiles;
-extern std::string_view ToolchainPrefix;
-} // namespace jit_compiler
+extern const resource_file ToolchainFiles[];
+extern unsigned long long NumToolchainFiles;
+extern resource_string_view ToolchainPrefix;
+} // namespace jit_compiler::resource

sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp

@@ -100,8 +100,12 @@ class HashPreprocessedAction : public PreprocessorFrontendAction {
 
 class SYCLToolchain {
   SYCLToolchain() {
+    using namespace jit_compiler::resource;
+
     for (size_t i = 0; i < NumToolchainFiles; ++i) {
-      auto [Path, Content] = ToolchainFiles[i];
+      resource_file RF = ToolchainFiles[i];
+      std::string_view Path{RF.Path.S, RF.Path.Size};
+      std::string_view Content{RF.Content.S, RF.Content.Size};
       ToolchainFS->addFile(Path, 0, llvm::MemoryBuffer::getMemBuffer(Content));
     }
   }
@@ -195,12 +199,14 @@ class SYCLToolchain {
     return std::move(Lib);
   }
 
+  std::string_view getPrefix() const { return Prefix; }
   std::string_view getClangXXExe() const { return ClangXXExe; }
 
 private:
   clang::IgnoringDiagConsumer IgnoreDiag;
-  std::string ClangXXExe =
-      (jit_compiler::ToolchainPrefix + "/bin/clang++").str();
+  std::string_view Prefix{jit_compiler::resource::ToolchainPrefix.S,
+                          jit_compiler::resource::ToolchainPrefix.Size};
+  std::string ClangXXExe = (Prefix + "/bin/clang++").str();
   llvm::IntrusiveRefCntPtr<llvm::vfs::InMemoryFileSystem> ToolchainFS =
       llvm::makeIntrusiveRefCnt<llvm::vfs::InMemoryFileSystem>();
 };
@@ -499,7 +505,7 @@ Error jit_compiler::linkDeviceLibraries(llvm::Module &Module,
   LLVMContext &Context = Module.getContext();
   for (const std::string &LibName : LibNames) {
     std::string LibPath =
-        (jit_compiler::ToolchainPrefix + "/lib/" + LibName).str();
+        (SYCLToolchain::instance().getPrefix() + "/lib/" + LibName).str();
 
     ModuleUPtr LibModule;
     if (auto Error = SYCLToolchain::instance()
@@ -518,7 +524,7 @@ Error jit_compiler::linkDeviceLibraries(llvm::Module &Module,
   // For GPU targets we need to link against vendor provided libdevice.
   if (IsCudaHIP) {
     Triple T{Module.getTargetTriple()};
-    Driver D{(jit_compiler::ToolchainPrefix + "/bin/clang++").str(),
+    Driver D{(SYCLToolchain::instance().getPrefix() + "/bin/clang++").str(),
              T.getTriple(), Diags};
     auto [CPU, Features] =
         Translator::getTargetCPUAndFeatureAttrs(&Module, "", Format);
@@ -764,7 +770,7 @@ jit_compiler::performPostLink(ModuleUPtr Module,
     auto &Ctx = Modules.front()->getContext();
     auto WrapLibraryInDevImg = [&](const std::string &LibName) -> Error {
       std::string LibPath =
-          (jit_compiler::ToolchainPrefix + "/lib/" + LibName).str();
+          (SYCLToolchain::instance().getPrefix() + "/lib/" + LibName).str();
       ModuleUPtr LibModule;
       if (auto Error = SYCLToolchain::instance()
                            .loadBitcodeLibrary(LibPath, Ctx)

sycl-jit/jit-compiler/utils/generate.py

@@ -28,8 +28,8 @@ def main():
             """
 #include <Resource.h>
 
-namespace jit_compiler {
-const std::pair<std::string_view, std::string_view> ToolchainFiles[] = {"""
+namespace jit_compiler::resource {
+const resource_file ToolchainFiles[] = {"""
         )
 
         def process_file(file_path):
@@ -41,7 +41,7 @@ def process_file(file_path):
     static const char data[] = {{
     #embed "{file_path}" if_empty(0)
         , 0}};
-    return std::string_view(data, std::size(data) - 1);
+    return resource_string_view{{data}};
     }}()
 }},"""
             )
@@ -66,9 +66,9 @@ def process_dir(dir):
             f"""
 }};
 
-size_t NumToolchainFiles = std::size(ToolchainFiles);
-std::string_view ToolchainPrefix = "{args.prefix}";
-}} // namespace jit_compiler
+unsigned long long NumToolchainFiles = size(ToolchainFiles);
+resource_string_view ToolchainPrefix{{"{args.prefix}"}};
+}} // namespace jit_compiler::resource
 """
         )
 

sycl/include/sycl/ext/oneapi/bindless_images.hpp

@@ -1036,15 +1036,19 @@ DataT fetch_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "HintT must always be a recognized standard type");
 
 #ifdef __SYCL_DEVICE_ONLY__
+  // Convert the raw handle to an image and use FETCH_UNSAMPLED_IMAGE since
+  // fetch_image should not use the sampler
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return FETCH_SAMPLED_IMAGE(
+    return FETCH_UNSAMPLED_IMAGE(
         DataT,
-        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyRead<coordSize>),
         coords);
   } else {
-    return sycl::bit_cast<DataT>(FETCH_SAMPLED_IMAGE(
+    return sycl::bit_cast<DataT>(FETCH_UNSAMPLED_IMAGE(
         HintT,
-        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyRead<coordSize>),
         coords));
   }
 #else

sycl/source/detail/device_impl.hpp

@@ -1580,16 +1580,17 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
           .value_or(0);
     }
     CASE(ext_oneapi_clock_sub_group) {
-      // Will be updated in a follow-up UR patch.
-      return false;
+      return get_info_impl_nocheck<UR_DEVICE_INFO_CLOCK_SUB_GROUP_SUPPORT_EXP>()
+          .value_or(0);
     }
     CASE(ext_oneapi_clock_work_group) {
-      // Will be updated in a follow-up UR patch.
-      return false;
+      return get_info_impl_nocheck<
+                 UR_DEVICE_INFO_CLOCK_WORK_GROUP_SUPPORT_EXP>()
+          .value_or(0);
     }
     CASE(ext_oneapi_clock_device) {
-      // Will be updated in a follow-up UR patch.
-      return false;
+      return get_info_impl_nocheck<UR_DEVICE_INFO_CLOCK_DEVICE_SUPPORT_EXP>()
+          .value_or(0);
     }
     else {
       return false; // This device aspect has not been implemented yet.

sycl/source/detail/ur_device_info_ret_types.inc

@@ -193,4 +193,7 @@ MAP(UR_DEVICE_INFO_MIPMAP_SUPPORT_EXP, ur_bool_t)
 MAP(UR_DEVICE_INFO_TIMESTAMP_RECORDING_SUPPORT_EXP, ur_bool_t)
 MAP(UR_DEVICE_INFO_KERNEL_LAUNCH_CAPABILITIES, ur_kernel_launch_properties_flags_t)
 MAP(UR_DEVICE_INFO_MEMORY_EXPORT_EXPORTABLE_DEVICE_MEM_EXP, ur_bool_t)
+MAP(UR_DEVICE_INFO_CLOCK_SUB_GROUP_SUPPORT_EXP, ur_bool_t)
+MAP(UR_DEVICE_INFO_CLOCK_WORK_GROUP_SUPPORT_EXP, ur_bool_t)
+MAP(UR_DEVICE_INFO_CLOCK_DEVICE_SUPPORT_EXP, ur_bool_t)
 // clang-format on

sycl/test-e2e/Basic/buffer/buffer_create.cpp

@@ -8,6 +8,11 @@
 // RUN: %{run} %t.out 2>&1 | FileCheck %s
 // UNSUPPORTED: ze_debug
 
+// L0v2 adapter doesn't optimize buffer creation based on device type yet
+// (integrated buffer implementation needs more work).
+// UNSUPPORTED: level_zero_v2_adapter
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/20121
+
 #include <iostream>
 #include <level_zero/ze_api.h>
 #include <sycl/detail/core.hpp>

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

@@ -11,9 +11,10 @@
 #include <sycl/ext/oneapi/bindless_images.hpp>
 #include <sycl/usm.hpp>
 
-class kernel_sampled_fetch;
+namespace {
 
-int main() {
+template <typename T, sycl::image_channel_type ChanType>
+static int testSampledImageFetch() {
 
   sycl::device dev;
   sycl::queue q(dev);
@@ -23,9 +24,9 @@ int main() {
   constexpr size_t width = 5;
   constexpr size_t height = 6;
   constexpr size_t N = width * height;
-  std::vector<sycl::vec<uint16_t, 4>> out(N);
-  std::vector<sycl::vec<uint16_t, 4>> expected(N);
-  std::vector<sycl::vec<uint16_t, 4>> dataIn(N);
+  std::vector<sycl::vec<T, 4>> out(N);
+  std::vector<sycl::vec<T, 4>> expected(N);
+  std::vector<sycl::vec<T, 4>> dataIn(N);
   for (int i = 0; i < width; i++) {
     for (int j = 0; j < height; j++) {
       auto index = i + (width * j);
@@ -43,8 +44,7 @@ int main() {
         sycl::filtering_mode::linear);
 
     // Extension: image descriptor
-    syclexp::image_descriptor desc({width, height}, 4,
-                                   sycl::image_channel_type::unsigned_int16);
+    syclexp::image_descriptor desc({width, height}, 4, ChanType);
     size_t pitch = 0;
 
     // Extension: returns the device pointer to USM allocated pitched memory
@@ -65,21 +65,20 @@ int main() {
 
     sycl::buffer buf(out.data(), sycl::range{height, width});
     q.submit([&](sycl::handler &cgh) {
-      auto outAcc = buf.get_access<sycl::access_mode::write>(
+      auto outAcc = buf.template get_access<sycl::access_mode::write>(
           cgh, sycl::range<2>{height, width});
 
-      cgh.parallel_for<kernel_sampled_fetch>(
-          sycl::nd_range<2>{{width, height}, {width, height}},
-          [=](sycl::nd_item<2> it) {
-            size_t dim0 = it.get_local_id(0);
-            size_t dim1 = it.get_local_id(1);
+      cgh.parallel_for(sycl::nd_range<2>{{width, height}, {width, height}},
+                       [=](sycl::nd_item<2> it) {
+                         size_t dim0 = it.get_local_id(0);
+                         size_t dim1 = it.get_local_id(1);
 
-            // Extension: fetch data from sampled image handle
-            auto px1 = syclexp::fetch_image<sycl::vec<uint16_t, 4>>(
-                imgHandle, sycl::int2(dim0, dim1));
+                         // Extension: fetch data from sampled image handle
+                         auto px1 = syclexp::fetch_image<sycl::vec<T, 4>>(
+                             imgHandle, sycl::int2(dim0, dim1));
 
-            outAcc[sycl::id<2>{dim1, dim0}] = px1;
-          });
+                         outAcc[sycl::id<2>{dim1, dim0}] = px1;
+                       });
     });
 
     q.wait_and_throw();
@@ -121,3 +120,23 @@ int main() {
   std::cout << "Test failed!" << std::endl;
   return 3;
 }
+
+} // namespace
+
+int main() {
+  if (int err =
+          testSampledImageFetch<uint16_t,
+                                sycl::image_channel_type::unsigned_int16>()) {
+    return err;
+  }
+  if (int err =
+          testSampledImageFetch<uint32_t,
+                                sycl::image_channel_type::unsigned_int32>()) {
+    return err;
+  }
+  if (int err =
+          testSampledImageFetch<float, sycl::image_channel_type::fp32>()) {
+    return err;
+  }
+  return 0;
+}