Merge branch 'sycl' into counter-based-4

.github/workflows/scorecard.yml

@@ -57,6 +57,6 @@ jobs:
 
       # Upload the results to GitHub's code scanning dashboard.
       - name: "Upload to code-scanning"
-        uses: github/codeql-action/upload-sarif@e2b3eafc8d227b0241d48be5f425d47c2d750a13 # v3.26.10
+        uses: github/codeql-action/upload-sarif@662472033e021d55d94146f66f6058822b0b39fd # v3.27.0
         with:
           sarif_file: results.sarif

.github/workflows/sycl-nightly.yml

@@ -198,7 +198,7 @@ jobs:
           echo "TAG=$(date +'%Y-%m-%d')-${GITHUB_SHA::7}" >> "$GITHUB_OUTPUT"
         fi
     - name: Upload binaries
-      uses: softprops/action-gh-release@c062e08bd532815e2082a85e87e3ef29c3e6d191
+      uses: softprops/action-gh-release@e7a8f85e1c67a31e6ed99a94b41bd0b71bbee6b8
       with:
         files: |
           sycl_linux.tar.gz

devops/cts_exclude_filter_L0_GPU

@@ -3,5 +3,3 @@ kernel_bundle
 marray
 # fix: https://github.com/KhronosGroup/SYCL-CTS/pull/964
 accessor_legacy
-# CMPLRLLVM-62822
-multi_ptr

devops/cts_exclude_filter_OCL_CPU

@@ -7,5 +7,3 @@ math_builtin_api
 hierarchical
 # fix: https://github.com/KhronosGroup/SYCL-CTS/pull/964
 accessor_legacy
-# CMPLRLLVM-62822
-multi_ptr

llvm/docs/requirements.txt

@@ -1,4 +1,4 @@
-sphinx==8.0.2
+sphinx==8.1.3
 docutils==0.21.2
 sphinx-markdown-tables==0.0.17
 recommonmark==0.7.1

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -1,7 +1,7 @@
-# commit 3d58884b4939d9bd095c917f8dd823ac8486684c
-# Merge: 6ade245e b0bd146a
+# commit fa8cc8ec16c1a2cf0926cc64026edc6a254ff0c2
+# Merge: 3d58884b 1984ceb1
 # Author: aarongreig <[email protected]>
 # Date:   Thu Oct 31 14:05:55 2024 +0000
 #     Merge pull request #2228 from nrspruit/copy_engine_refactor
 #     [L0] Refactor Copy Engine Usage checks for Performance
-set(UNIFIED_RUNTIME_TAG 65ccdbd4355e4bc7b3e3e0f6d841a2d900871f0a)
+set(UNIFIED_RUNTIME_TAG 65ccdbd4355e4bc7b3e3e0f6d841a2d900871f0a)

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -10,26 +10,20 @@
 //  1. Add a new enumerator to
 //     `sycl::ext::oneapi::experimental::detail::PropKind` representing the new
 //     property. Increment
-//     `sycl::ext::oneapi::experimental::detail::PropKind::PropKindSize`
-//  2. Define property key class with `value_t` that must be `property_value`
-//     with the first template argument being the property class itself. The
-//     name of the key class must be the property name suffixed by `_key`, i.e.
-//     for a property `foo` the class should be named `foo_key`.
+//     `sycl::ext::oneapi::experimental::detail::PropKind::PropKindSize`.
+//  2. Define property key class inherited from
+//     `detail::compile_time_property_key` with `value_t` that must be
+//     `property_value` with the first template argument being the property
+//     class itself. The name of the key class must be the property name
+//     suffixed by `_key`, i.e. for a property `foo` the class should be named
+//     `foo_key`.
 //  3. Add an `inline constexpr` variable in the same namespace as the property
 //     key. The variable should have the same type as `value_t` of the property
 //     class, e.g. for a property `foo`, there should be a definition
 //     `inline constexpr foo_key::value_t foo`.
-//  4. Specialize `sycl::ext::oneapi::experimental::is_property_key` and
-//     `sycl::ext::oneapi::experimental::is_property_key_of` for the property
-//     key class.
-//  5. Specialize `sycl::ext::oneapi::experimental::detail::PropertyToKind` for
-//     the new property key class. The specialization should have a `Kind`
-//     member with the value equal to the enumerator added in 1.
-//  6. Specialize
-//     `sycl::ext::oneapi::experimental::detail::IsCompileTimeProperty` for the
-//     new property key class. This specialization should derive from
-//     `std::true_type`.
-//  7. If the property needs an LLVM IR attribute, specialize
+//  4. Specialize `sycl::ext::oneapi::experimental::is_property_key_of` for the
+//     property key class.
+//  5. If the property needs an LLVM IR attribute, specialize
 //     `sycl::ext::oneapi::experimental::detail::PropertyMetaInfo` for the new
 //     `value_t` of the property key class. The specialization must have a
 //     `static constexpr const char *name` member with a value equal to the

sycl/include/sycl/multi_ptr.hpp

@@ -732,12 +732,28 @@ class __SYCL_TYPE(multi_ptr) multi_ptr<void, Space, DecorateAddress> {
   decorated_type *m_Pointer;
 };
 
+namespace detail {
+// See access.hpp's DecoratedType<..., access::address_space::constant_space>.
+//
+// This is only applicable to `access::decorated::legacy` mode because constant
+// AS is deprecated itself and is only accessible in legacy modes.
+template <auto Space>
+#ifdef __SYCL_DEVICE_ONLY__
+inline constexpr auto decoration_space =
+    deduce_AS<typename DecoratedType<void, Space>::type>::value;
+#else
+inline constexpr auto decoration_space = Space;
+#endif
+} // namespace detail
+
 // Legacy specialization of multi_ptr.
 // TODO: Add deprecation warning here when possible.
 template <typename ElementType, access::address_space Space>
 class __SYCL2020_DEPRECATED(
     "decorated::legacy multi_ptr specialization is deprecated since SYCL 2020.")
     multi_ptr<ElementType, Space, access::decorated::legacy> {
+  static constexpr auto DecorationSpace = detail::decoration_space<Space>;
+
 public:
   using value_type = ElementType;
   using element_type =
@@ -777,7 +793,8 @@ class __SYCL2020_DEPRECATED(
 
   multi_ptr(ElementType *pointer)
       : m_Pointer(detail::dynamic_address_cast<
-                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {
+                  DecorationSpace, /* SupressNotImplementedAssert = */ true>(
+            pointer)) {
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
   }
@@ -786,7 +803,8 @@ class __SYCL2020_DEPRECATED(
   template <typename = typename detail::const_if_const_AS<Space, ElementType>>
   multi_ptr(const ElementType *pointer)
       : m_Pointer(detail::dynamic_address_cast<
-                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {}
+                  DecorationSpace, /* SupressNotImplementedAssert = */ true>(
+            pointer)) {}
 #endif
 
   multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
@@ -814,7 +832,7 @@ class __SYCL2020_DEPRECATED(
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
     m_Pointer = detail::dynamic_address_cast<
-        Space, /* SupressNotImplementedAssert = */ true>(pointer);
+        DecorationSpace, /* SupressNotImplementedAssert = */ true>(pointer);
     return *this;
   }
 
@@ -856,8 +874,8 @@ class __SYCL2020_DEPRECATED(
   multi_ptr(accessor<ElementType, dimensions, Mode, target::device,
                      isPlaceholder, PropertyListT>
                 Accessor)
-      : multi_ptr(
-            detail::static_address_cast<Space>(Accessor.get_pointer().get())) {}
+      : multi_ptr(detail::static_address_cast<DecorationSpace>(
+            Accessor.get_pointer().get())) {}
 
   // Only if Space == local_space || generic_space
   template <
@@ -1088,6 +1106,8 @@ template <access::address_space Space>
 class __SYCL2020_DEPRECATED(
     "decorated::legacy multi_ptr specialization is deprecated since SYCL 2020.")
     multi_ptr<void, Space, access::decorated::legacy> {
+  static constexpr auto DecorationSpace = detail::decoration_space<Space>;
+
 public:
   using value_type = void;
   using element_type = void;
@@ -1113,17 +1133,17 @@ class __SYCL2020_DEPRECATED(
                                   !std::is_same_v<RelayPointerT, void *>>>
   multi_ptr(void *pointer)
       : m_Pointer(detail::dynamic_address_cast<
-                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {
+                  DecorationSpace, /* SupressNotImplementedAssert = */ true>(
+            pointer)) {
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
   }
 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
   template <typename = typename detail::const_if_const_AS<Space, void>>
   multi_ptr(const void *pointer)
-      : m_Pointer(
-            detail::dynamic_address_cast<
-                pointer_t, /* SupressNotImplementedAssert = */ true>(pointer)) {
-  }
+      : m_Pointer(detail::dynamic_address_cast<
+                  DecorationSpace, /* SupressNotImplementedAssert = */ true>(
+            pointer)) {}
 #endif
 #endif
   multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
@@ -1154,7 +1174,7 @@ class __SYCL2020_DEPRECATED(
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
     m_Pointer = detail::dynamic_address_cast<
-        Space, /* SupressNotImplementedAssert = */ true>(pointer);
+        DecorationSpace, /* SupressNotImplementedAssert = */ true>(pointer);
     return *this;
   }
 #endif
@@ -1249,6 +1269,8 @@ template <access::address_space Space>
 class __SYCL2020_DEPRECATED(
     "decorated::legacy multi_ptr specialization is deprecated since SYCL 2020.")
     multi_ptr<const void, Space, access::decorated::legacy> {
+  static constexpr auto DecorationSpace = detail::decoration_space<Space>;
+
 public:
   using value_type = const void;
   using element_type = const void;
@@ -1275,15 +1297,17 @@ class __SYCL2020_DEPRECATED(
                                   !std::is_same_v<RelayPointerT, const void *>>>
   multi_ptr(const void *pointer)
       : m_Pointer(detail::dynamic_address_cast<
-                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {
+                  DecorationSpace, /* SupressNotImplementedAssert = */ true>(
+            pointer)) {
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
   }
 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
   template <typename = typename detail::const_if_const_AS<Space, void>>
   multi_ptr(const void *pointer)
       : m_Pointer(detail::dynamic_address_cast<
-                  Space, /* SupressNotImplementedAssert = */ true>(pointer)) {}
+                  DecorationSpace, /* SupressNotImplementedAssert = */ true>(
+            pointer)) {}
 #endif
 #endif
   multi_ptr(std::nullptr_t) : m_Pointer(nullptr) {}
@@ -1314,7 +1338,7 @@ class __SYCL2020_DEPRECATED(
     // TODO An implementation should reject an argument if the deduced
     // address space is not compatible with Space.
     m_Pointer = detail::dynamic_address_cast<
-        pointer_t, /* SupressNotImplementedAssert = */ true>(pointer);
+        DecorationSpace, /* SupressNotImplementedAssert = */ true>(pointer);
     return *this;
   }
 #endif
@@ -1442,7 +1466,7 @@ address_space_cast(ElementType *pointer) {
   // space is not compatible with Space.
   // Use LegacyPointerTypes here to also allow constant_space
   return multi_ptr<ElementType, Space, DecorateAddress>(
-      detail::dynamic_address_cast<Space,
+      detail::dynamic_address_cast<detail::decoration_space<Space>,
                                    /* SupressNotImplementedAssert = */ true>(
           pointer));
 }

sycl/source/detail/device_impl.cpp

@@ -675,7 +675,9 @@ bool device_impl::has(aspect Aspect) const {
     const std::vector<arch> supported_archs = {
         arch::intel_cpu_spr,     arch::intel_cpu_gnr,
         arch::intel_gpu_pvc,     arch::intel_gpu_dg2_g10,
-        arch::intel_gpu_dg2_g11, arch::intel_gpu_dg2_g12};
+        arch::intel_gpu_dg2_g11, arch::intel_gpu_dg2_g12,
+        arch::intel_gpu_bmg_g21, arch::intel_gpu_lnl_m,
+        arch::intel_gpu_arl_h};
     try {
       return std::any_of(
           supported_archs.begin(), supported_archs.end(),

sycl/source/detail/kernel_program_cache.hpp

@@ -9,6 +9,7 @@
 #pragma once
 
 #include "sycl/exception.hpp"
+#include <detail/config.hpp>
 #include <detail/kernel_arg_mask.hpp>
 #include <detail/platform_impl.hpp>
 #include <sycl/detail/common.hpp>
@@ -19,8 +20,10 @@
 
 #include <atomic>
 #include <condition_variable>
+#include <iomanip>
 #include <mutex>
 #include <set>
+#include <thread>
 #include <type_traits>
 
 #include <boost/unordered/unordered_flat_map.hpp>
@@ -176,6 +179,42 @@ class KernelProgramCache {
 
   void setContextPtr(const ContextPtr &AContext) { MParentContext = AContext; }
 
+  // Sends message to std:cerr stream when SYCL_CACHE_TRACE environemnt is
+  // set.
+  static inline void traceProgram(const std::string &Msg,
+                                  const ProgramCacheKeyT &CacheKey) {
+    if (!SYCLConfig<SYCL_CACHE_TRACE>::isTraceInMemCache())
+      return;
+
+    int ImageId = CacheKey.first.second;
+    std::stringstream DeviceList;
+    for (const auto &Device : CacheKey.second)
+      DeviceList << "0x" << std::setbase(16)
+                 << reinterpret_cast<uintptr_t>(Device) << ",";
+
+    std::string Identifier = "[Key:{imageId = " + std::to_string(ImageId) +
+                             ",urDevice = " + DeviceList.str() + "}]: ";
+
+    std::cerr << "[In-Memory Cache][Thread Id:" << std::this_thread::get_id()
+              << "][Program Cache]" << Identifier << Msg << std::endl;
+  }
+
+  // Sends message to std:cerr stream when SYCL_CACHE_TRACE environemnt is
+  // set.
+  static inline void traceKernel(const std::string &Msg,
+                                 const std::string &KernelName,
+                                 bool IsKernelFastCache = false) {
+    if (!SYCLConfig<SYCL_CACHE_TRACE>::isTraceInMemCache())
+      return;
+
+    std::string Identifier =
+        "[IsFastCache: " + std::to_string(IsKernelFastCache) +
+        "][Key:{Name = " + KernelName + "}]: ";
+
+    std::cerr << "[In-Memory Cache][Thread Id:" << std::this_thread::get_id()
+              << "][Kernel Cache]" << Identifier << Msg << std::endl;
+  }
+
   Locked<ProgramCache> acquireCachedPrograms() {
     return {MCachedPrograms, MProgramCacheMutex};
   }
@@ -195,7 +234,9 @@ class KernelProgramCache {
       CommonProgramKeyT CommonKey =
           std::make_pair(CacheKey.first.second, CacheKey.second);
       ProgCache.KeyMap.emplace(CommonKey, CacheKey);
-    }
+      traceProgram("Program inserted.", CacheKey);
+    } else
+      traceProgram("Program fetched.", CacheKey);
     return std::make_pair(It->second, DidInsert);
   }
 
@@ -217,7 +258,9 @@ class KernelProgramCache {
       CommonProgramKeyT CommonKey =
           std::make_pair(CacheKey.first.second, CacheKey.second);
       ProgCache.KeyMap.emplace(CommonKey, CacheKey);
-    }
+      traceProgram("Program inserted.", CacheKey);
+    } else
+      traceProgram("Program fetched.", CacheKey);
     return DidInsert;
   }
 
@@ -227,8 +270,11 @@ class KernelProgramCache {
     auto LockedCache = acquireKernelsPerProgramCache();
     auto &Cache = LockedCache.get()[Program];
     auto [It, DidInsert] = Cache.try_emplace(KernelName, nullptr);
-    if (DidInsert)
+    if (DidInsert) {
       It->second = std::make_shared<KernelBuildResult>(getAdapter());
+      traceKernel("Kernel inserted.", KernelName);
+    } else
+      traceKernel("Kernel fetched.", KernelName);
     return std::make_pair(It->second, DidInsert);
   }
 
@@ -237,6 +283,7 @@ class KernelProgramCache {
     std::unique_lock<std::mutex> Lock(MKernelFastCacheMutex);
     auto It = MKernelFastCache.find(CacheKey);
     if (It != MKernelFastCache.end()) {
+      traceKernel("Kernel fetched.", std::get<3>(CacheKey), true);
       return It->second;
     }
     return std::make_tuple(nullptr, nullptr, nullptr, nullptr);
@@ -247,6 +294,7 @@ class KernelProgramCache {
     std::unique_lock<std::mutex> Lock(MKernelFastCacheMutex);
     // if no insertion took place, thus some other thread has already inserted
     // smth in the cache
+    traceKernel("Kernel inserted.", std::get<3>(CacheKey), true);
     MKernelFastCache.emplace(CacheKey, CacheVal);
   }
 

sycl/test-e2e/Basic/built-ins.cpp

@@ -4,8 +4,9 @@
 // RUN: %{build} -D__SYCL_USE_VARIADIC_SPIRV_OCL_PRINTF__ -Wno-#warnings -o %t_var.out
 // RUN: %{run} %t_var.out | FileCheck %s
 
-// Hits an assertion with AMD:
-// XFAIL: hip_amd
+// Hits an assertion and kernel page fault with AMD:
+// UNSUPPORTED: hip_amd
+// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/14404
 
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/experimental/builtins.hpp>