Merge remote-tracking branch 'intel/origin/sycl' into maronas/ext_com…

…posite_device

devops/dependencies.json

@@ -1,15 +1,15 @@
 {
   "linux": {
     "compute_runtime": {
-      "github_tag": "23.48.27912.11",
-      "version": "23.48.27912.11",
-      "url": "https://github.com/intel/compute-runtime/releases/tag/23.48.27912.11",
+      "github_tag": "23.52.28202.14",
+      "version": "23.52.28202.14",
+      "url": "https://github.com/intel/compute-runtime/releases/tag/23.52.28202.14",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     },
     "igc": {
-      "github_tag": "igc-1.0.15610.11",
-      "version": "1.0.15610.11",
-      "url": "https://github.com/intel/intel-graphics-compiler/releases/tag/igc-1.0.15610.11",
+      "github_tag": "igc-1.0.15770.11",
+      "version": "1.0.15770.11",
+      "url": "https://github.com/intel/intel-graphics-compiler/releases/tag/igc-1.0.15770.11",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     },
     "cm": {

sycl/doc/extensions/proposed/sycl_ext_oneapi_forward_progress.asciidoc

@@ -20,7 +20,7 @@
 == Notice
 
 [%hardbreaks]
-Copyright (C) 2022-2023 Intel Corporation.  All rights reserved.
+Copyright (C) 2022-2024 Intel Corporation.  All rights reserved.
 
 Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
 of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
@@ -383,9 +383,11 @@ and work-groups to also provide concurrent forward progress guarantees). In
 such a case, an implementation must satisfy the strongest request(s).
 
 Devices may not be able to provide the requested forward progress guarantees
-for all launch configurations.  The <<launch, launch queries>> defined in a
-later section allow developers to identify valid launch configurations for
-specific combinations of properties.
+for all launch configurations. Developers should use the launch queries defined
+by the
+link:../proposed/sycl_ext_oneapi_launch_queries.asciidoc[sycl_ext_oneapi_launch_queries]
+extension to identify valid launch configurations for specific combinations of
+properties.
 
 [NOTE]
 ====

sycl/doc/extensions/proposed/sycl_ext_oneapi_root_group.asciidoc

@@ -182,7 +182,8 @@ inline constexpr use_root_sync_key::value_t use_root_sync;
 === The `root_group` class
 
 The `root_group` class implements all member functions common to the
-`sycl::group` and `sycl::sub_group` classes.
+`sycl::group` and `sycl::sub_group` classes and also contains own
+additional functions.
 
 [source,c++]
 ----
@@ -191,6 +192,13 @@ namespace ext {
 namespace oneapi {
 namespace experimental {
 
+enum class execution_scope {
+  work_item,
+  sub_group,
+  work_group,
+  root_group,
+};
+
 template <int Dimensions>
 class root_group {
 public:
@@ -221,6 +229,31 @@ public:
 
   bool leader() const;
 
+  template <execution_scope Scope>
+  std::enable_if_t<(Scope == execution_scope::work_item ||
+                    Scope == execution_scope::work_group),
+                   id<Dimensions>>
+  get_id() const;
+
+  template <execution_scope Scope>
+  std::enable_if_t<Scope == execution_scope::sub_group, id<1>> get_id() const;
+
+  template <execution_scope Scope>
+  size_t get_linear_id() const;
+
+  template <execution_scope Scope>
+  std::enable_if_t<(Scope == execution_scope::work_item ||
+                    Scope == execution_scope::work_group),
+                   range<Dimensions>>
+  get_range() const;
+
+  template <execution_scope Scope>
+  std::enable_if_t<Scope == execution_scope::sub_group, range<1>>
+  get_range() const;
+
+  template <execution_scope Scope>
+  size_t get_linear_range() const;
+
 };
 
 } // namespace experimental
@@ -307,6 +340,67 @@ work-item is the leader of the root-group, and `false` for all other work-items
 in the root-group. The leader of the root-group is guaranteed to be the
 work-item for which `get_local_id()` returns 0.
 
+[source,c++]
+----
+template <execution_scope Scope>
+std::enable_if_t<(Scope == execution_scope::work_item ||
+                  Scope == execution_scope::work_group),
+                 id<Dimensions>>
+get_id() const;
+----
+_Returns_: An `id` representing the index of the current work-group or work-item at `Scope`
+hierarchy level within the `root_group` object.
+
+[source,c++]
+----
+template <execution_scope Scope>
+std::enable_if_t<Scope == execution_scope::sub_group, id<1>> get_id() const;
+----
+_Returns_: An `id` representing the index of the current sub-group within the
+`root_group` object.
+
+[source,c++]
+----
+template <execution_scope Scope>
+size_t get_linear_id() const;
+----
+_Constraints_: `Scope` must be narrower than
+`execution_scope::root_group`.
+
+_Returns_: A linearized number of the current work-group or work-item at `Scope` hierarchy
+level within the `root_group` object.
+
+[source,c++]
+----
+template <execution_scope Scope>
+std::enable_if_t<(Scope == execution_scope::work_item ||
+                  Scope == execution_scope::work_group),
+                 range<Dimensions>>
+get_range() const;
+----
+_Returns_: A `range` representing the number of work-groups or work-items of `Scope`
+hierarchy level within the `root_group` object.
+
+[source,c++]
+----
+template <execution_scope Scope>
+std::enable_if_t<Scope == execution_scope::sub_group, range<1>>
+get_range() const;
+----
+_Returns_: A `range` representing the number of sub-groups within the `root_group`
+object.
+
+[source,c++]
+----
+template <execution_scope Scope>
+size_t get_linear_range() const;
+----
+_Constraints_: `Scope` must be narrower than
+`execution_scope::root_group`.
+
+_Returns_: The number of work-groups or work-items of `Scope` hierarchy level within the
+`root_group` object.
+
 
 === Using a `root_group`
 

sycl/include/sycl/detail/builtins/common_functions.inc

@@ -72,32 +72,15 @@ min(T x, detail::get_elem_type_t<T> y) {
              detail::simplify_if_swizzle_t<T>{y});
 }
 
-#undef BUILTIN_COMMON
-
-#ifdef __SYCL_DEVICE_ONLY__
-DEVICE_IMPL_TEMPLATE(THREE_ARGS, clamp, builtin_enable_generic_t,
-                     [](auto... xs) {
-                       using ElemTy = detail::get_elem_type_t<T0>;
-                       if constexpr (std::is_integral_v<ElemTy>) {
-                         if constexpr (std::is_signed_v<ElemTy>) {
-                           return __spirv_ocl_s_clamp(xs...);
-                         } else {
-                           return __spirv_ocl_u_clamp(xs...);
-                         }
-                       } else {
-                         return __spirv_ocl_fclamp(xs...);
-                       }
-                     })
-#else
-HOST_IMPL_TEMPLATE(THREE_ARGS, clamp, builtin_enable_generic_t, common,
-                   default_ret_type)
-#endif
+BUILTIN_COMMON(THREE_ARGS, clamp, __spirv_ocl_fclamp)
 template <typename T>
-detail::builtin_enable_generic_non_scalar_t<T>
+detail::builtin_enable_common_non_scalar_t<T>
 clamp(T x, detail::get_elem_type_t<T> y, detail::get_elem_type_t<T> z) {
   return clamp(detail::simplify_if_swizzle_t<T>{x},
                detail::simplify_if_swizzle_t<T>{y},
                detail::simplify_if_swizzle_t<T>{z});
 }
+
+#undef BUILTIN_COMMON
 } // namespace _V1
 } // namespace sycl

sycl/include/sycl/detail/builtins/integer_functions.inc

@@ -122,6 +122,15 @@ min(T x, detail::get_elem_type_t<T> y) {
              detail::simplify_if_swizzle_t<T>{y});
 }
 
+BUILTIN_GENINT_SU(THREE_ARGS, clamp)
+template <typename T>
+detail::builtin_enable_integer_non_scalar_t<T>
+clamp(T x, detail::get_elem_type_t<T> y, detail::get_elem_type_t<T> z) {
+  return clamp(detail::simplify_if_swizzle_t<T>{x},
+               detail::simplify_if_swizzle_t<T>{y},
+               detail::simplify_if_swizzle_t<T>{z});
+}
+
 BUILTIN_GENINT(ONE_ARG, clz)
 BUILTIN_GENINT(ONE_ARG, ctz)
 BUILTIN_GENINT(ONE_ARG, popcount)