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

  CONFLICT (content): Merge conflict in sycl/source/detail/graph_memory_pool.cpp

Author: mikolaj-pirog

.github/CODEOWNERS

@@ -165,6 +165,8 @@ libdevice/nativecpu* @intel/dpcpp-nativecpu-reviewers
 sycl/include/sycl/ext/oneapi/experimental/graph.hpp @intel/sycl-graphs-reviewers
 sycl/source/detail/graph_impl.cpp @intel/sycl-graphs-reviewers
 sycl/source/detail/graph_impl.hpp @intel/sycl-graphs-reviewers
+sycl/source/detail/graph_memory_pool.hpp @intel/sycl-graphs-reviewers
+sycl/source/detail/graph_memory_pool.cpp @intel/sycl-graphs-reviewers
 sycl/unittests/Extensions/CommandGraph/ @intel/sycl-graphs-reviewers
 sycl/test-e2e/Graph @intel/sycl-graphs-reviewers
 sycl/doc/design/CommandGraph.md @intel/sycl-graphs-reviewers

.github/workflows/sycl-nightly.yml

@@ -217,6 +217,9 @@ jobs:
   run-sycl-cts-linux:
     needs: [ubuntu2204_build, build-sycl-cts-linux]
     if: ${{ always() && !cancelled() && needs.ubuntu2204_build.outputs.build_conclusion == 'success' }}
+    # Normally these jobs take less than 10m. But sometimes the job hangs up and
+    # reaches the 360m limit. Set a lower limit to free up the runner earlier.
+    timeout-minutes: 35
     strategy:
       fail-fast: false
       matrix:

libdevice/sanitizer/tsan_rtl.cpp

@@ -134,40 +134,9 @@ inline __SYCL_GLOBAL__ RawShadow *MemToShadow(uptr addr, uint32_t as) {
   return shadow_ptr;
 }
 
-// For CPU device, each work group is a thread
-inline Sid GetCurrentSid_CPU() {
-  // work-group linear id
-  const auto wg_lid =
-      __spirv_BuiltInWorkgroupId.x * __spirv_BuiltInNumWorkgroups.y *
-          __spirv_BuiltInNumWorkgroups.z +
-      __spirv_BuiltInWorkgroupId.y * __spirv_BuiltInNumWorkgroups.z +
-      __spirv_BuiltInWorkgroupId.z;
-  return wg_lid;
-}
-
-// For GPU device, each work item is a thread
-inline Sid GetCurrentSid_GPU() {
-  // sub-group linear id
-  const auto lid = __spirv_BuiltInGlobalLinearId;
-  return lid;
-}
-
 inline Sid GetCurrentSid() {
-#if defined(__LIBDEVICE_CPU__)
-  return GetCurrentSid_CPU();
-#elif defined(__LIBDEVICE_PVC__)
-  return GetCurrentSid_GPU();
-#else
-  if (TsanLaunchInfo->DeviceTy == DeviceType::CPU) {
-    return GetCurrentSid_CPU();
-  } else if (TsanLaunchInfo->DeviceTy != DeviceType::UNKNOWN) {
-    return GetCurrentSid_GPU();
-  } else {
-    TSAN_DEBUG(__spirv_ocl_printf(__tsan_print_unsupport_device_type,
-                                  (int)TsanLaunchInfo->DeviceTy));
-    return 0;
-  }
-#endif
+  const auto lid = __spirv_BuiltInGlobalLinearId;
+  return lid % kThreadSlotCount;
 }
 
 inline RawShadow LoadShadow(const __SYCL_GLOBAL__ RawShadow *p) {
@@ -466,10 +435,6 @@ DEVICE_EXTERN_C_NOINLINE void __tsan_cleanup_private(uptr addr, uint32_t size) {
 
 DEVICE_EXTERN_C_INLINE void __tsan_device_barrier() {
   Sid sid = GetCurrentSid();
-  __spirv_ControlBarrier(__spv::Scope::Device, __spv::Scope::Device,
-                         __spv::MemorySemanticsMask::SequentiallyConsistent |
-                             __spv::MemorySemanticsMask::CrossWorkgroupMemory |
-                             __spv::MemorySemanticsMask::WorkgroupMemory);
 
   // sync current thread clock to global state
   TsanLaunchInfo->Clock[kThreadSlotCount].clk_[sid] =
@@ -484,19 +449,10 @@ DEVICE_EXTERN_C_INLINE void __tsan_device_barrier() {
   for (uptr i = 0; i < kThreadSlotCount; i++)
     TsanLaunchInfo->Clock[sid].clk_[i] =
         TsanLaunchInfo->Clock[kThreadSlotCount].clk_[i];
-
-  __spirv_ControlBarrier(__spv::Scope::Device, __spv::Scope::Device,
-                         __spv::MemorySemanticsMask::SequentiallyConsistent |
-                             __spv::MemorySemanticsMask::CrossWorkgroupMemory |
-                             __spv::MemorySemanticsMask::WorkgroupMemory);
 }
 
-static inline void __tsan_group_barrier_impl() {
+DEVICE_EXTERN_C_INLINE void __tsan_group_barrier() {
   Sid sid = GetCurrentSid();
-  __spirv_ControlBarrier(__spv::Scope::Workgroup, __spv::Scope::Workgroup,
-                         __spv::MemorySemanticsMask::SequentiallyConsistent |
-                             __spv::MemorySemanticsMask::CrossWorkgroupMemory |
-                             __spv::MemorySemanticsMask::WorkgroupMemory);
 
   // sync current thread clock to global state
   TsanLaunchInfo->Clock[kThreadSlotCount].clk_[sid] =
@@ -511,23 +467,6 @@ static inline void __tsan_group_barrier_impl() {
   for (uptr i = 0; i < kThreadSlotCount; i++)
     TsanLaunchInfo->Clock[sid].clk_[i] =
         TsanLaunchInfo->Clock[kThreadSlotCount].clk_[i];
-
-  __spirv_ControlBarrier(__spv::Scope::Workgroup, __spv::Scope::Workgroup,
-                         __spv::MemorySemanticsMask::SequentiallyConsistent |
-                             __spv::MemorySemanticsMask::CrossWorkgroupMemory |
-                             __spv::MemorySemanticsMask::WorkgroupMemory);
-}
-
-DEVICE_EXTERN_C_INLINE void __tsan_group_barrier() {
-#if defined(__LIBDEVICE_CPU__)
-  return;
-#elif defined(__LIBDEVICE_PVC__)
-  __tsan_group_barrier_impl();
-#else
-  if (TsanLaunchInfo->DeviceTy == DeviceType::CPU)
-    return;
-  __tsan_group_barrier_impl();
-#endif
 }
 
 #endif // __SPIR__ || __SPIRV__

llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp

@@ -361,8 +361,7 @@ void ThreadSanitizerOnSpirv::appendDebugInfoToArgs(
   auto &Loc = I->getDebugLoc();
 
   // SPIR constant address space
-  PointerType *ConstASPtrTy =
-      PointerType::get(Type::getInt8Ty(C), kSpirOffloadConstantAS);
+  PointerType *ConstASPtrTy = PointerType::get(C, kSpirOffloadConstantAS);
 
   // File & Line
   if (Loc) {

sycl/include/sycl/detail/util.hpp

@@ -90,10 +90,10 @@ template <> struct ABINeutralT<std::vector<std::string>> {
 template <typename T> using ABINeutralT_t = typename ABINeutralT<T>::type;
 
 template <typename ParamT> auto convert_to_abi_neutral(ParamT &&Info) {
-  using ParamNoRef = std::remove_reference_t<ParamT>;
-  if constexpr (std::is_same_v<ParamNoRef, std::string>) {
+  using ParamDecayT = std::decay_t<ParamT>;
+  if constexpr (std::is_same_v<ParamDecayT, std::string>) {
     return detail::string{Info};
-  } else if constexpr (std::is_same_v<ParamNoRef, std::vector<std::string>>) {
+  } else if constexpr (std::is_same_v<ParamDecayT, std::vector<std::string>>) {
     std::vector<detail::string> Res;
     Res.reserve(Info.size());
     for (std::string &Str : Info) {

sycl/source/detail/device_impl.hpp

@@ -471,7 +471,7 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
 
 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
   template <typename Param, bool InitializingCache = false>
-  typename Param::return_type get_info() const {
+  decltype(auto) get_info() const {
 #define CALL_GET_INFO get_info
 #else
   // We've been exporting
@@ -484,10 +484,27 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
 #define CALL_GET_INFO get_info_abi_workaround
   template <typename Param> typename Param::return_type get_info() const;
   template <typename Param, bool InitializingCache = false>
-  typename Param::return_type get_info_abi_workaround() const {
+  decltype(auto) get_info_abi_workaround() const {
 #endif
     using execution_scope = ext::oneapi::experimental::execution_scope;
 
+    // With the return type of this function being automatically
+    // deduced we can't simply do
+    //
+    //    CASE(Desc1) { return get_info<Desc2>(); }
+    //
+    // because the function isn't defined yet and we can't auto-deduce the
+    // return type for `Desc2` yet. The solution here is to make that delegation
+    // template-parameter-dependent. We use the `InitializingCache` parameter
+    // for that out of convenience.
+    //
+    // Note that for "eager" cache it's the programmer's responsibility that
+    // the descriptor we delegate to is initialized first (by referencing that
+    // descriptor first when defining the cache data member). For "CallOnce"
+    // cache we want to be querying cached value so "false" is the right
+    // template parameter for such delegation.
+    constexpr bool DependentFalse = InitializingCache && false;
+
     if constexpr (decltype(MCache)::has<Param>() && !InitializingCache) {
       return MCache.get<Param>();
     }
@@ -523,11 +540,13 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
       return range<3>{result[2], result[1], result[0]};
     }
     CASE(info::device::max_work_item_sizes<2>) {
-      range<3> r3 = CALL_GET_INFO<info::device::max_work_item_sizes<3>>();
+      range<3> r3 =
+          CALL_GET_INFO<info::device::max_work_item_sizes<3>, DependentFalse>();
       return range<2>{r3[1], r3[2]};
     }
     CASE(info::device::max_work_item_sizes<1>) {
-      range<3> r3 = CALL_GET_INFO<info::device::max_work_item_sizes<3>>();
+      range<3> r3 =
+          CALL_GET_INFO<info::device::max_work_item_sizes<3>, DependentFalse>();
       return range<1>{r3[2]};
     }
 
@@ -608,16 +627,18 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
       // profiling, urDeviceGetGlobalTimestamps is not supported,
       // command_submit, command_start, command_end will be calculated. See
       // MFallbackProfiling
-      return get_info_impl<UR_DEVICE_INFO_QUEUE_PROPERTIES>() &
-             UR_QUEUE_FLAG_PROFILING_ENABLE;
+      return static_cast<bool>(
+          get_info_impl<UR_DEVICE_INFO_QUEUE_PROPERTIES>() &
+          UR_QUEUE_FLAG_PROFILING_ENABLE);
     }
 
     CASE(info::device::built_in_kernels) {
       return split_string(get_info_impl<UR_DEVICE_INFO_BUILT_IN_KERNELS>(),
                           ';');
     }
     CASE(info::device::built_in_kernel_ids) {
-      auto names = CALL_GET_INFO<info::device::built_in_kernels>();
+      auto names =
+          CALL_GET_INFO<info::device::built_in_kernels, DependentFalse>();
 
       std::vector<kernel_id> ids;
       ids.reserve(names.size());
@@ -655,7 +676,8 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
             "the info::device::preferred_interop_user_sync info descriptor can "
             "only be queried with an OpenCL backend");
 
-      return get_info_impl<UR_DEVICE_INFO_PREFERRED_INTEROP_USER_SYNC>();
+      return static_cast<bool>(
+          get_info_impl<UR_DEVICE_INFO_PREFERRED_INTEROP_USER_SYNC>());
     }
 
     CASE(info::device::partition_properties) {
@@ -753,16 +775,19 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
     }
 
     CASE(info::device::usm_device_allocations) {
-      return get_info_impl<UR_DEVICE_INFO_USM_DEVICE_SUPPORT>() &
-             UR_DEVICE_USM_ACCESS_CAPABILITY_FLAG_ACCESS;
+      return static_cast<bool>(
+          get_info_impl<UR_DEVICE_INFO_USM_DEVICE_SUPPORT>() &
+          UR_DEVICE_USM_ACCESS_CAPABILITY_FLAG_ACCESS);
     }
     CASE(info::device::usm_host_allocations) {
-      return get_info_impl<UR_DEVICE_INFO_USM_HOST_SUPPORT>() &
-             UR_DEVICE_USM_ACCESS_CAPABILITY_FLAG_ACCESS;
+      return static_cast<bool>(
+          get_info_impl<UR_DEVICE_INFO_USM_HOST_SUPPORT>() &
+          UR_DEVICE_USM_ACCESS_CAPABILITY_FLAG_ACCESS);
     }
     CASE(info::device::usm_shared_allocations) {
-      return get_info_impl<UR_DEVICE_INFO_USM_SINGLE_SHARED_SUPPORT>() &
-             UR_DEVICE_USM_ACCESS_CAPABILITY_FLAG_ACCESS;
+      return static_cast<bool>(
+          get_info_impl<UR_DEVICE_INFO_USM_SINGLE_SHARED_SUPPORT>() &
+          UR_DEVICE_USM_ACCESS_CAPABILITY_FLAG_ACCESS);
     }
     CASE(info::device::usm_restricted_shared_allocations) {
       ur_device_usm_access_capability_flags_t cap_flags =
@@ -773,14 +798,16 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
                 UR_DEVICE_USM_ACCESS_CAPABILITY_FLAG_CONCURRENT_ACCESS));
     }
     CASE(info::device::usm_system_allocations) {
-      return get_info_impl<UR_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT>() &
-             UR_DEVICE_USM_ACCESS_CAPABILITY_FLAG_ACCESS;
+      return static_cast<bool>(
+          get_info_impl<UR_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT>() &
+          UR_DEVICE_USM_ACCESS_CAPABILITY_FLAG_ACCESS);
     }
 
     CASE(info::device::opencl_c_version) {
       throw sycl::exception(errc::feature_not_supported,
                             "Deprecated interface that hasn't been working for "
                             "some time already");
+      return std::string{}; // for return type deduction.
     }
 
     CASE(ext::intel::info::device::max_mem_bandwidth) {
@@ -794,30 +821,34 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
     CASE(info::device::ext_oneapi_max_global_work_groups) {
       // Deprecated alias.
       return CALL_GET_INFO<
-          ext::oneapi::experimental::info::device::max_global_work_groups>();
+          ext::oneapi::experimental::info::device::max_global_work_groups,
+          DependentFalse>();
     }
     CASE(info::device::ext_oneapi_max_work_groups_1d) {
       // Deprecated alias.
       return CALL_GET_INFO<
-          ext::oneapi::experimental::info::device::max_work_groups<1>>();
+          ext::oneapi::experimental::info::device::max_work_groups<1>,
+          DependentFalse>();
     }
     CASE(info::device::ext_oneapi_max_work_groups_2d) {
       // Deprecated alias.
       return CALL_GET_INFO<
-          ext::oneapi::experimental::info::device::max_work_groups<2>>();
+          ext::oneapi::experimental::info::device::max_work_groups<2>,
+          DependentFalse>();
     }
     CASE(info::device::ext_oneapi_max_work_groups_3d) {
       // Deprecated alias.
       return CALL_GET_INFO<
-          ext::oneapi::experimental::info::device::max_work_groups<3>>();
+          ext::oneapi::experimental::info::device::max_work_groups<3>,
+          DependentFalse>();
     }
 
     CASE(info::device::ext_oneapi_cuda_cluster_group) {
       if (getBackend() != backend::ext_oneapi_cuda)
         return false;
 
       return get_info_impl_nocheck<UR_DEVICE_INFO_CLUSTER_LAUNCH_SUPPORT_EXP>()
-          .value_or(0);
+                 .value_or(0) != 0;
     }
 
     // ext_codeplay_device_traits.def
@@ -834,7 +865,8 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
     }
     CASE(ext::oneapi::experimental::info::device::max_work_groups<3>) {
       size_t Limit = CALL_GET_INFO<
-          ext::oneapi::experimental::info::device::max_global_work_groups>();
+          ext::oneapi::experimental::info::device::max_global_work_groups,
+          DependentFalse>();
 
       // TODO: std::array<size_t, 3> ?
       size_t result[3];
@@ -846,12 +878,14 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
     }
     CASE(ext::oneapi::experimental::info::device::max_work_groups<2>) {
       id<3> max_3d = CALL_GET_INFO<
-          ext::oneapi::experimental::info::device::max_work_groups<3>>();
+          ext::oneapi::experimental::info::device::max_work_groups<3>,
+          DependentFalse>();
       return id<2>{max_3d[1], max_3d[2]};
     }
     CASE(ext::oneapi::experimental::info::device::max_work_groups<1>) {
       id<3> max_3d = CALL_GET_INFO<
-          ext::oneapi::experimental::info::device::max_work_groups<3>>();
+          ext::oneapi::experimental::info::device::max_work_groups<3>,
+          DependentFalse>();
       return id<1>{max_3d[2]};
     }
 
@@ -896,8 +930,8 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
     }
 
     CASE(ext::oneapi::experimental::info::device::mipmap_max_anisotropy) {
-      // Implicit conversion:
-      return get_info_impl<UR_DEVICE_INFO_MIPMAP_MAX_ANISOTROPY_EXP>();
+      return static_cast<float>(
+          get_info_impl<UR_DEVICE_INFO_MIPMAP_MAX_ANISOTROPY_EXP>());
     }
 
     CASE(ext::oneapi::experimental::info::device::component_devices) {
@@ -906,7 +940,7 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
       if (!Devs.has_val()) {
         ur_result_t Err = Devs.error();
         if (Err == UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION)
-          return {};
+          return std::vector<sycl::device>{};
         getAdapter()->checkUrResult(Err);
       }
 
@@ -935,8 +969,8 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
                             "must have a composite device.");
     }
     CASE(ext::oneapi::info::device::num_compute_units) {
-      // uint32_t -> size_t
-      return get_info_impl<UR_DEVICE_INFO_NUM_COMPUTE_UNITS>();
+      return static_cast<size_t>(
+          get_info_impl<UR_DEVICE_INFO_NUM_COMPUTE_UNITS>());
     }
 
     // ext_intel_device_traits.def
@@ -1028,8 +1062,8 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
       return get_info_impl<UR_DEVICE_INFO_MEMORY_BUS_WIDTH>();
     }
     CASE(ext::intel::info::device::max_compute_queue_indices) {
-      // uint32_t->int implicit conversion.
-      return get_info_impl<UR_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES>();
+      return static_cast<int>(
+          get_info_impl<UR_DEVICE_INFO_MAX_COMPUTE_QUEUE_INDICES>());
     }
     CASE(ext::intel::esimd::info::device::has_2d_block_io_support) {
       if (!has(aspect::ext_intel_esimd))
@@ -1093,7 +1127,7 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
     }
     else {
       constexpr auto Desc = UrInfoCode<Param>::value;
-      return get_info_impl<Desc>();
+      return static_cast<typename Param::return_type>(get_info_impl<Desc>());
     }
 #undef CASE
   }