[llvm] [libsycl] Implement nd_range kernel submissions (PR #206505)
Kseniya Tikhomirova via llvm-commits
llvm-commits at lists.llvm.org
Thu Jul 16 03:11:17 PDT 2026
https://github.com/KseniyaTikhomirova updated https://github.com/llvm/llvm-project/pull/206505
>From da69f104e7d5bba0f294f5c540211cdfbdef5f0b Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Fri, 29 May 2026 08:08:12 -0700
Subject: [PATCH 1/5] nd_range impl
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/docs/index.rst | 2 +-
.../sycl/__impl/detail/kernel_arg_helpers.hpp | 13 ++
.../sycl/__impl/detail/linearization.hpp | 42 ++++
.../sycl/__impl/detail/unified_range_view.hpp | 5 +
libsycl/include/sycl/__impl/group.hpp | 147 ++++++++++++++
libsycl/include/sycl/__impl/memory_enums.hpp | 37 ++++
libsycl/include/sycl/__impl/nd_item.hpp | 183 ++++++++++++++++++
libsycl/include/sycl/__impl/nd_range.hpp | 89 +++++++++
libsycl/include/sycl/__impl/queue.hpp | 70 +++++--
libsycl/include/sycl/__impl/sub_group.hpp | 112 +++++++++++
libsycl/include/sycl/__spirv/spirv_vars.hpp | 14 ++
libsycl/include/sycl/sycl.hpp | 5 +
libsycl/src/detail/queue_impl.cpp | 12 ++
libsycl/test/basic/group.cpp | 78 ++++++++
libsycl/test/basic/group_local_id.cpp | 47 +++++
libsycl/test/basic/linear_sub_group.cpp | 37 ++++
libsycl/test/basic/nd_range.cpp | 54 ++++++
.../test/basic/queue_parallel_for_generic.cpp | 24 ++-
.../basic/sub_group_by_value_semantics.cpp | 39 ++++
libsycl/test/basic/sub_group_common.cpp | 80 ++++++++
20 files changed, 1073 insertions(+), 17 deletions(-)
create mode 100644 libsycl/include/sycl/__impl/detail/linearization.hpp
create mode 100644 libsycl/include/sycl/__impl/group.hpp
create mode 100644 libsycl/include/sycl/__impl/memory_enums.hpp
create mode 100644 libsycl/include/sycl/__impl/nd_item.hpp
create mode 100644 libsycl/include/sycl/__impl/nd_range.hpp
create mode 100644 libsycl/include/sycl/__impl/sub_group.hpp
create mode 100644 libsycl/test/basic/group.cpp
create mode 100644 libsycl/test/basic/group_local_id.cpp
create mode 100644 libsycl/test/basic/linear_sub_group.cpp
create mode 100644 libsycl/test/basic/nd_range.cpp
create mode 100644 libsycl/test/basic/sub_group_by_value_semantics.cpp
create mode 100644 libsycl/test/basic/sub_group_common.cpp
diff --git a/libsycl/docs/index.rst b/libsycl/docs/index.rst
index a5d657f3e6404..233d5682fe9c8 100644
--- a/libsycl/docs/index.rst
+++ b/libsycl/docs/index.rst
@@ -113,7 +113,7 @@ TODO for added SYCL classes
* to implement submit & copy with accessors (low priority)
* get_info & properties
* ctors that accepts context (blocked by lack of liboffload support)
- * nd_range kernel submissions
+ * nd_range kernel submissions: offset is not supported by liboffload now, SYCL2020 deprecated feature
* cross-context events wait (host tasks are needed)
* implement check if lambda arguments are device copyable (requires clang support of corresponding builtins) unless FE will fully cover it
* kernel instantiating on host (debugging purposes)
diff --git a/libsycl/include/sycl/__impl/detail/kernel_arg_helpers.hpp b/libsycl/include/sycl/__impl/detail/kernel_arg_helpers.hpp
index 31a50e835cc80..4967f8657c28e 100644
--- a/libsycl/include/sycl/__impl/detail/kernel_arg_helpers.hpp
+++ b/libsycl/include/sycl/__impl/detail/kernel_arg_helpers.hpp
@@ -15,6 +15,7 @@
#define _LIBSYCL___IMPL_DETAIL_KERNEL_ARG_HELPERS
#include <sycl/__impl/index_space_classes.hpp>
+#include <sycl/__impl/nd_item.hpp>
#include <sycl/__impl/detail/config.hpp>
@@ -126,6 +127,18 @@ class Builder {
return __spirv::initBuiltInGlobalInvocationId<Dims, id<Dims>>();
}
+ /// \return the work group currently being operated on by the device.
+ template <int Dims> static const group<Dims> getElement(group<Dims> *) {
+ static_assert(isValidDimensions<Dims>, "invalid dimensions");
+ return group<Dims>();
+ }
+
+ /// \return the nd_item currently being operated on by the device.
+ template <int Dims> static const nd_item<Dims> getElement(nd_item<Dims> *) {
+ static_assert(isValidDimensions<Dims>, "invalid dimensions");
+ return nd_item<Dims>();
+ }
+
/// Constructs item with the given data.
/// \param Extent a range representing the dimensions of the range of possible
/// values of the item.
diff --git a/libsycl/include/sycl/__impl/detail/linearization.hpp b/libsycl/include/sycl/__impl/detail/linearization.hpp
new file mode 100644
index 0000000000000..c1ad4151f6aa8
--- /dev/null
+++ b/libsycl/include/sycl/__impl/detail/linearization.hpp
@@ -0,0 +1,42 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Helpers for SYCL index/range linearization. Follows SYCL2020 3.11.1.
+/// Linearization.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_DETAIL_LINEARIZATION_HPP
+#define _LIBSYCL___IMPL_DETAIL_LINEARIZATION_HPP
+
+#include <sycl/__impl/index_space_classes.hpp>
+
+#include <cstddef>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+
+template <int Dimensions>
+inline std::size_t linearize_id(const id<Dimensions> &Index,
+ const range<Dimensions> &Extent) noexcept {
+ if constexpr (Dimensions == 1) {
+ return Index[0];
+ } else if constexpr (Dimensions == 2) {
+ return Index[0] * Extent[1] + Index[1];
+ } else {
+ return Index[0] * Extent[1] * Extent[2] + Index[1] * Extent[2] + Index[2];
+ }
+}
+
+} // namespace detail
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_DETAIL_LINEARIZATION_HPP
diff --git a/libsycl/include/sycl/__impl/detail/unified_range_view.hpp b/libsycl/include/sycl/__impl/detail/unified_range_view.hpp
index 464ce6524df2c..be53d0e6331b6 100644
--- a/libsycl/include/sycl/__impl/detail/unified_range_view.hpp
+++ b/libsycl/include/sycl/__impl/detail/unified_range_view.hpp
@@ -37,6 +37,11 @@ struct UnifiedRangeView {
UnifiedRangeView(sycl::range<Dims> &N)
: MGlobalSize(&(N[0])), MDims(size_t(Dims)) {}
+ template <int Dims>
+ UnifiedRangeView(sycl::nd_range<Dims> &N)
+ : MGlobalSize(&(N.MGlobalSize[0])), MLocalSize(&(N.MLocalSize[0])),
+ MOffset(&(N.MOffset[0])), MDims{size_t(Dims)} {}
+
UnifiedRangeView(const size_t *GlobalSize, const size_t *LocalSize,
const size_t *Offset, size_t Dims)
: MGlobalSize(GlobalSize), MLocalSize(LocalSize), MOffset(Offset),
diff --git a/libsycl/include/sycl/__impl/group.hpp b/libsycl/include/sycl/__impl/group.hpp
new file mode 100644
index 0000000000000..edadcbf4cbacb
--- /dev/null
+++ b/libsycl/include/sycl/__impl/group.hpp
@@ -0,0 +1,147 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains the declaration of the SYCL 2020 group class.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_GROUP_HPP
+#define _LIBSYCL___IMPL_GROUP_HPP
+
+#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/detail/linearization.hpp>
+#include <sycl/__impl/index_space_classes.hpp>
+#include <sycl/__impl/memory_enums.hpp>
+#include <sycl/__spirv/spirv_vars.hpp>
+
+#include <cstddef>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+template <int> class nd_item;
+
+// SYCL2020 4.9.1.7. group class.
+/// The group class encapsulates all functionality required to represent a
+/// particular work-group within a parallel execution.
+template <int Dimensions = 1> class group {
+public:
+ using id_type = id<Dimensions>;
+ using range_type = range<Dimensions>;
+ using linear_id_type = std::size_t;
+ static constexpr int dimensions = Dimensions;
+ static constexpr memory_scope fence_scope = memory_scope::work_group;
+
+ group(const group &rhs) = default;
+ group(group &&rhs) = default;
+ group &operator=(const group &rhs) = default;
+ group &operator=(group &&rhs) = default;
+
+ /// \return an id representing the index of the work-group within the global
+ /// nd-range for every dimension.
+ id<Dimensions> get_group_id() const noexcept {
+ return __spirv::initBuiltInWorkgroupId<Dimensions, id<Dimensions>>();
+ }
+
+ /// Equivalent to `return get_group_id()[dimension]`.
+ std::size_t get_group_id(int dimension) const noexcept {
+ return get_group_id()[dimension];
+ }
+
+ /// \return a SYCL id representing the calling work-item’s position within the
+ /// work-group.
+ id<Dimensions> get_local_id() const noexcept {
+ return __spirv::initBuiltInLocalInvocationId<Dimensions, id<Dimensions>>();
+ }
+
+ /// Equivalent to `return get_local_id()[dimension]`.
+ std::size_t get_local_id(int dimension) const noexcept {
+ return get_local_id()[dimension];
+ }
+
+ /// \return a SYCL range representing all dimensions of the local range.
+ range<Dimensions> get_local_range() const noexcept {
+ return __spirv::initBuiltInWorkgroupSize<Dimensions, range<Dimensions>>();
+ }
+
+ /// Equivalent to `return get_local_range()[dimension]`.
+ std::size_t get_local_range(int dimension) const noexcept {
+ return get_local_range()[dimension];
+ }
+
+ /// \return a SYCL range representing the number of work-groups in the
+ /// nd-range.
+ range<Dimensions> get_group_range() const noexcept {
+ return __spirv::initBuiltInNumWorkgroups<Dimensions, range<Dimensions>>();
+ }
+
+ /// Equivalent to `return get_group_range()[dimension]`.
+ std::size_t get_group_range(int dimension) const noexcept {
+ return get_group_range()[dimension];
+ }
+
+ /// \return a SYCL range representing the maximum number of work-items in any
+ /// work-group in the nd-range.
+ range<Dimensions> get_max_local_range() const noexcept {
+ return get_local_range();
+ }
+
+ /// Equivalent to `return get_group_id(dimension)`.
+ std::size_t operator[](int dimension) const noexcept {
+ return get_group_id(dimension);
+ }
+
+ /// \return the linearized work-group id within the nd-range.
+ std::size_t get_group_linear_id() const noexcept {
+ return detail::linearize_id(get_group_id(), get_group_range());
+ }
+
+ /// \return a linearized version of the calling work-item’s local id.
+ std::size_t get_local_linear_id() const noexcept {
+ return detail::linearize_id(get_local_id(), get_local_range());
+ }
+
+ /// \return the total number of work-groups in the nd-range.
+ std::size_t get_group_linear_range() const noexcept {
+ auto groupRange = get_group_range();
+ return multiply_all_dims(groupRange);
+ }
+
+ /// \return the total number of work-items in this work-group.
+ std::size_t get_local_linear_range() const noexcept {
+ auto localRange = get_local_range();
+ return multiply_all_dims(localRange);
+ }
+
+ /// \return true for exactly one work-item in the work-group, if the calling
+ /// work-item is the leader of the work-group, and false for all other
+ /// work-items in the work-group.
+ bool leader() const noexcept { return (get_local_linear_id() == 0); }
+
+ // TODO: implement parallel_for_work_item, async_work_group_copy and wait_for.
+
+protected:
+ group() = default;
+
+ static std::size_t
+ multiply_all_dims(const range<Dimensions> &Range) noexcept {
+ if constexpr (Dimensions == 1) {
+ return Range[0];
+ } else if constexpr (Dimensions == 2) {
+ return Range[0] * Range[1];
+ } else {
+ return Range[0] * Range[1] * Range[2];
+ }
+ }
+
+ template <int dimensions> friend class sycl::nd_item;
+};
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_GROUP_HPP
diff --git a/libsycl/include/sycl/__impl/memory_enums.hpp b/libsycl/include/sycl/__impl/memory_enums.hpp
new file mode 100644
index 0000000000000..37fbb1111465d
--- /dev/null
+++ b/libsycl/include/sycl/__impl/memory_enums.hpp
@@ -0,0 +1,37 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains SYCL 2020 memory scope enumeration (3.8.3.2.).
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_MEMORY_ENUMS_HPP
+#define _LIBSYCL___IMPL_MEMORY_ENUMS_HPP
+
+#include <sycl/__impl/detail/config.hpp>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+enum class memory_scope : std::uint32_t {
+ work_item = 0,
+ sub_group = 1,
+ work_group = 2,
+ device = 3,
+ system = 4
+};
+
+inline constexpr auto memory_scope_work_item = memory_scope::work_item;
+inline constexpr auto memory_scope_sub_group = memory_scope::sub_group;
+inline constexpr auto memory_scope_work_group = memory_scope::work_group;
+inline constexpr auto memory_scope_device = memory_scope::device;
+inline constexpr auto memory_scope_system = memory_scope::system;
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_MEMORY_ENUMS_HPP
diff --git a/libsycl/include/sycl/__impl/nd_item.hpp b/libsycl/include/sycl/__impl/nd_item.hpp
new file mode 100644
index 0000000000000..632269349c925
--- /dev/null
+++ b/libsycl/include/sycl/__impl/nd_item.hpp
@@ -0,0 +1,183 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains the declaration of the SYCL 2020 nd_item class.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_ND_ITEM_HPP
+#define _LIBSYCL___IMPL_ND_ITEM_HPP
+
+#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/detail/linearization.hpp>
+#include <sycl/__impl/group.hpp>
+#include <sycl/__impl/index_space_classes.hpp>
+#include <sycl/__impl/nd_range.hpp>
+#include <sycl/__impl/sub_group.hpp>
+#include <sycl/__spirv/spirv_vars.hpp>
+
+#include <cstddef>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+class Builder;
+} // namespace detail
+
+// SYCL2020 4.9.1.5. nd_item class.
+/// nd_item<int Dimensions> identifies an instance of the function object
+/// executing at each point in an nd_range<int Dimensions> passed to a
+/// parallel_for call.
+template <int Dimensions = 1> class nd_item {
+public:
+ static constexpr int dimensions = Dimensions;
+
+ nd_item(const nd_item &rhs) = default;
+ nd_item(nd_item &&rhs) = default;
+ nd_item &operator=(const nd_item &rhs) = default;
+ nd_item &operator=(nd_item &&rhs) = default;
+
+ friend bool operator==(const nd_item &lhs, const nd_item &rhs) {
+ // https://github.com/KhronosGroup/SYCL-Docs/issues/532
+ return true;
+ }
+
+ friend bool operator!=(const nd_item &lhs, const nd_item &rhs) {
+ return !(lhs == rhs);
+ }
+
+ /// \return the constituent global id representing the work-item’s position in
+ /// the global iteration space.
+ id<Dimensions> get_global_id() const noexcept {
+ return __spirv::initBuiltInGlobalInvocationId<Dimensions, id<Dimensions>>();
+ }
+
+ /// \return the constituent element of the global id representing the
+ /// work-item’s position in the nd-range in the given Dimension.
+ std::size_t get_global_id(int dimension) const noexcept {
+ return get_global_id()[dimension];
+ }
+
+ /// \return the constituent global id as a linear index value, representing
+ /// the work-item’s position in the global iteration space.
+ std::size_t get_global_linear_id() const noexcept {
+ id<Dimensions> adjustedIndex = get_global_id();
+ const id<Dimensions> offset =
+ __spirv::initBuiltInGlobalOffset<Dimensions, id<Dimensions>>();
+ if constexpr (Dimensions == 1) {
+ adjustedIndex[0] -= offset[0];
+ } else if constexpr (Dimensions == 2) {
+ adjustedIndex[0] -= offset[0];
+ adjustedIndex[1] -= offset[1];
+ } else {
+ adjustedIndex[0] -= offset[0];
+ adjustedIndex[1] -= offset[1];
+ adjustedIndex[2] -= offset[2];
+ }
+ return detail::linearize_id(adjustedIndex, get_global_range());
+ }
+
+ /// \return the constituent local id representing the work-item’s position
+ /// within the current work-group.
+ id<Dimensions> get_local_id() const noexcept {
+ return __spirv::initBuiltInLocalInvocationId<Dimensions, id<Dimensions>>();
+ }
+
+ /// \return the constituent element of the local id representing the
+ /// work-item’s position within the current work-group in the given Dimension.
+ std::size_t get_local_id(int dimension) const noexcept {
+ return get_local_id()[dimension];
+ }
+
+ /// \return the constituent local id as a linear index value, representing the
+ /// work-item’s position within the current work-group.
+ std::size_t get_local_linear_id() const noexcept {
+ return detail::linearize_id(get_local_id(), get_local_range());
+ }
+
+ /// \return the constituent work-group, group representing the work-group's
+ /// position within the overall nd-range.
+ group<Dimensions> get_group() const noexcept { return group<Dimensions>(); }
+
+ /// \return a sub_group representing the sub-group to which the work-item
+ /// belongs.
+ sub_group get_sub_group() const noexcept { return sub_group(); }
+
+ /// \return the constituent element of the group id representing the
+ /// work-group’s position within the overall nd_range in the given Dimension.
+ std::size_t get_group(int dimension) const noexcept {
+ return get_group_id()[dimension];
+ }
+
+ /// \return the group id as a linear index value.
+ std::size_t get_group_linear_id() const noexcept {
+ return detail::linearize_id(get_group_id(), get_group_range());
+ }
+
+ /// \return the number of work-groups in the iteration space.
+ range<Dimensions> get_group_range() const noexcept {
+ return __spirv::initBuiltInNumWorkgroups<Dimensions, range<Dimensions>>();
+ }
+
+ /// \return the number of work-groups for Dimension in the iteration space.
+ std::size_t get_group_range(int dimension) const noexcept {
+ return get_group_range()[dimension];
+ }
+
+ /// \return a range representing the dimensions of the global iteration space.
+ range<Dimensions> get_global_range() const noexcept {
+ return __spirv::initBuiltInGlobalSize<Dimensions, range<Dimensions>>();
+ }
+
+ /// Equivalent to return get_global_range().get(dimension).
+ std::size_t get_global_range(int dimension) const noexcept {
+ return get_global_range()[dimension];
+ }
+
+ /// \return a range representing the dimensions of the current work-group.
+ range<Dimensions> get_local_range() const noexcept {
+ return __spirv::initBuiltInWorkgroupSize<Dimensions, range<Dimensions>>();
+ }
+
+ /// Equivalent to return get_local_range().get(dimension).
+ std::size_t get_local_range(int dimension) const noexcept {
+ return get_local_range()[dimension];
+ }
+
+ /// Deprecated in SYCL 2020.
+ /// \return an id representing the n-dimensional offset provided to the
+ /// constructor of the nd_range and that is added by the runtime to the global
+ /// id of each work-item.
+ __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
+ id<Dimensions> get_offset() const noexcept {
+ return __spirv::initBuiltInGlobalOffset<Dimensions, id<Dimensions>>();
+ }
+
+ /// \return the nd_range of the current execution.
+ nd_range<Dimensions> get_nd_range() const noexcept {
+ return nd_range<Dimensions>(
+ get_global_range(), get_local_range(),
+ __spirv::initBuiltInGlobalOffset<Dimensions, id<Dimensions>>());
+ }
+
+ // TODO: add wait_for and async_work_group_copy once builtins are implemented.
+
+protected:
+ friend class detail::Builder;
+
+ nd_item() = default;
+
+ id<Dimensions> get_group_id() const {
+ return __spirv::initBuiltInWorkgroupId<Dimensions, id<Dimensions>>();
+ }
+};
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_ND_ITEM_HPP
diff --git a/libsycl/include/sycl/__impl/nd_range.hpp b/libsycl/include/sycl/__impl/nd_range.hpp
new file mode 100644
index 0000000000000..b3d292e1b2e36
--- /dev/null
+++ b/libsycl/include/sycl/__impl/nd_range.hpp
@@ -0,0 +1,89 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains the declaration of the SYCL 2020 nd_range class.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_ND_RANGE_HPP
+#define _LIBSYCL___IMPL_ND_RANGE_HPP
+
+#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/index_space_classes.hpp>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+class UnifiedRangeView;
+} // namespace detail
+
+// SYCL 2020 4.9.1.2. nd_range class.
+/// nd_range<int Dimensions> defines the iteration domain of both the
+/// work-groups and the overall dispatch.
+template <int Dimensions = 1> class nd_range {
+ static_assert(Dimensions >= 1 && Dimensions <= 3,
+ "nd_range can only be 1-, 2-, or 3-dimensional.");
+
+public:
+ static constexpr int dimensions = Dimensions;
+
+ nd_range(const nd_range<Dimensions> &rhs) = default;
+ nd_range(nd_range<Dimensions> &&rhs) = default;
+ nd_range<Dimensions> &operator=(const nd_range<Dimensions> &rhs) = default;
+ nd_range<Dimensions> &operator=(nd_range<Dimensions> &&rhs) = default;
+
+ friend bool operator==(const nd_range<Dimensions> &lhs,
+ const nd_range<Dimensions> &rhs) {
+ return (rhs.MGlobalSize == lhs.MGlobalSize) &&
+ (rhs.MLocalSize == lhs.MLocalSize) && (rhs.MOffset == lhs.MOffset);
+ }
+
+ friend bool operator!=(const nd_range<Dimensions> &lhs,
+ const nd_range<Dimensions> &rhs) {
+ return !(lhs == rhs);
+ }
+
+ __SYCL2020_DEPRECATED("offset is deprecated in SYCL2020")
+ nd_range(range<Dimensions> globalSize, range<Dimensions> localSize,
+ id<Dimensions> offset) noexcept
+ : MGlobalSize(globalSize), MLocalSize(localSize), MOffset(offset) {}
+
+ nd_range(range<Dimensions> globalSize, range<Dimensions> localSize)
+ : MGlobalSize(globalSize), MLocalSize(localSize),
+ MOffset(id<Dimensions>()) {}
+
+ /// \return the constituent global range.
+ range<Dimensions> get_global_range() const noexcept { return MGlobalSize; }
+
+ /// \return the constituent local range.
+ range<Dimensions> get_local_range() const noexcept { return MLocalSize; }
+
+ /// This range would result from globalSize/localSize as provided on
+ /// construction.
+ /// \return a range representing the number of groups in each dimension.
+ range<Dimensions> get_group_range() const noexcept {
+ return MGlobalSize / MLocalSize;
+ }
+
+ /// Deprecated in SYCL 2020.
+ /// \return the constituent offset.
+ __SYCL2020_DEPRECATED("offset is deprecated in SYCL2020")
+ id<Dimensions> get_offset() const noexcept { return MOffset; }
+
+protected:
+ range<Dimensions> MGlobalSize;
+ range<Dimensions> MLocalSize;
+ id<Dimensions> MOffset;
+
+ friend class detail::UnifiedRangeView;
+};
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_ND_RANGE_HPP
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index d05592e5fa3fd..4a2a8dd0ed4e0 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -332,9 +332,37 @@ class _LIBSYCL_EXPORT queue {
std::forward<Rest>(rest)...);
}
+ template <typename KernelName = detail::AutoName, int Dims, typename... Rest>
+ event parallel_for(nd_range<Dims> executionRange, Rest &&...rest) {
+ return parallel_for<KernelName, Dims, Rest...>(executionRange, {},
+ std::forward<Rest>(rest)...);
+ }
+
+ template <typename KernelName = detail::AutoName, int Dims, typename... Rest>
+ event parallel_for(nd_range<Dims> executionRange, event depEvent,
+ Rest &&...rest) {
+ return parallel_for<KernelName, Dims, Rest...>(executionRange, {depEvent},
+ std::forward<Rest>(rest)...);
+ }
+
+ template <typename KernelName = detail::AutoName, int Dims, typename... Rest>
+ event parallel_for(nd_range<Dims> executionRange,
+ const std::vector<event> &depEvents, Rest &&...rest) {
+ if (executionRange.get_global_range() != range<Dims>{} &&
+ (executionRange.get_local_range() == range<Dims>{} ||
+ executionRange.get_global_range() % executionRange.get_local_range() !=
+ range<Dims>{}))
+ throw sycl::exception(sycl::make_error_code(sycl::errc::nd_range),
+ "Invalid nd_range submission: global size must be "
+ "evenly divisible by local size.");
+ return parallelForImpl<KernelName>(executionRange, depEvents,
+ std::forward<Rest>(rest)...);
+ }
+
private:
- template <typename KernelName, int Dims, typename... Rest>
- event parallelForImpl(range<Dims> numWorkItems,
+ template <typename KernelName, int Dims, template <int> class Range,
+ typename... Rest>
+ event parallelForImpl(Range<Dims> numWorkItems,
const std::vector<event> &depEvents, Rest &&...rest) {
if constexpr (sizeof...(Rest) != 1)
throw sycl::exception(errc::feature_not_supported,
@@ -343,21 +371,37 @@ class _LIBSYCL_EXPORT queue {
using KernelType =
std::decay_t<detail::nth_type_t<sizeof...(Rest) - 1, Rest...>>;
- using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
- static_assert(
- std::is_convertible_v<sycl::item<Dims>, LambdaArgType> ||
- std::is_convertible_v<sycl::item<Dims, false>, LambdaArgType>,
- "Kernel argument of a sycl::parallel_for with sycl::range "
- "must be either sycl::item or be convertible from sycl::item");
- using TranformedLambdaArgType = std::conditional_t<
- std::is_convertible_v<item<Dims>, LambdaArgType>, item<Dims>,
+ constexpr bool IsNdRangeSubmission =
+ std::is_same_v<Range<Dims>, nd_range<Dims>>;
+ using SuggestedArgType =
+ std::conditional_t<IsNdRangeSubmission, nd_item<Dims>, item<Dims>>;
+ using LambdaArgType =
+ sycl::detail::lambda_arg_type<KernelType, SuggestedArgType>;
+
+ if constexpr (IsNdRangeSubmission) {
+ static_assert(
+ std::is_convertible_v<sycl::nd_item<Dims>, LambdaArgType>,
+ "Kernel argument of a sycl::parallel_for with sycl::nd_range "
+ "must be sycl::nd_item");
+ } else {
+ static_assert(
+ std::is_convertible_v<sycl::item<Dims>, LambdaArgType> ||
+ std::is_convertible_v<sycl::item<Dims, false>, LambdaArgType>,
+ "Kernel argument of a sycl::parallel_for with sycl::range "
+ "must be sycl::item or be convertible from sycl::item");
+ }
+
+ using TransformedLambdaArgType = std::conditional_t<
+ IsNdRangeSubmission, nd_item<Dims>,
std::conditional_t<
- std::is_convertible_v<item<Dims, false>, LambdaArgType>,
- item<Dims, false>, LambdaArgType>>;
+ std::is_convertible_v<sycl::item<Dims>, LambdaArgType>, item<Dims>,
+ std::conditional_t<
+ std::is_convertible_v<sycl::item<Dims, false>, LambdaArgType>,
+ item<Dims, false>, LambdaArgType>>>;
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
- submitParallelFor<NameT, TranformedLambdaArgType, KernelType>(rest...);
+ submitParallelFor<NameT, TransformedLambdaArgType, KernelType>(rest...);
return getLastEvent();
}
diff --git a/libsycl/include/sycl/__impl/sub_group.hpp b/libsycl/include/sycl/__impl/sub_group.hpp
new file mode 100644
index 0000000000000..fbcba75170e03
--- /dev/null
+++ b/libsycl/include/sycl/__impl/sub_group.hpp
@@ -0,0 +1,112 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains the declaration of the SYCL 2020 sub_group class.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_SUB_GROUP_HPP
+#define _LIBSYCL___IMPL_SUB_GROUP_HPP
+
+#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/index_space_classes.hpp>
+#include <sycl/__impl/memory_enums.hpp>
+#include <sycl/__spirv/spirv_vars.hpp>
+
+#include <cstdint>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+template <int> class nd_item;
+
+// SYCL 2020 4.9.1.8. sub_group class.
+/// The sub_group class encapsulates all functionality required to represent a
+/// particular sub-group within a parallel execution.
+class sub_group {
+public:
+ using id_type = id<1>;
+ using range_type = sycl::range<1>;
+ using linear_id_type = std::uint32_t;
+ static constexpr int dimensions = 1;
+ static constexpr memory_scope fence_scope = memory_scope::sub_group;
+
+ sub_group(const sub_group &rhs) = default;
+ sub_group(sub_group &&rhs) = default;
+ sub_group &operator=(const sub_group &rhs) = default;
+ sub_group &operator=(sub_group &&rhs) = default;
+
+ friend bool operator==(const sub_group &lhs, const sub_group &rhs) {
+ return lhs.get_group_id() == rhs.get_group_id();
+ }
+
+ friend bool operator!=(const sub_group &lhs, const sub_group &rhs) {
+ return !(lhs == rhs);
+ }
+
+ /// \return an id representing the index of the sub-group within the
+ /// work-group.
+ id_type get_group_id() const noexcept { return __spirv_BuiltInSubgroupId(); }
+
+ /// \return a SYCL id representing the calling work-item’s position within the
+ /// sub-group.
+ id_type get_local_id() const noexcept {
+ return __spirv_BuiltInSubgroupLocalInvocationId();
+ }
+
+ /// \return a range representing the size of the sub-group.
+ range_type get_local_range() const noexcept {
+ return __spirv_BuiltInSubgroupSize();
+ }
+
+ /// \return a range representing the number of sub-groups within the
+ /// work-group.
+ range_type get_group_range() const noexcept {
+ return __spirv_BuiltInNumSubgroups();
+ }
+
+ /// \return a range representing the maximum number of work-items permitted in
+ /// a sub-group for the executing kernel.
+ range_type get_max_local_range() const noexcept {
+ return __spirv_BuiltInSubgroupMaxSize();
+ }
+
+ /// Equivalent to return get_group_id()[0].
+ linear_id_type get_group_linear_id() const noexcept {
+ return static_cast<linear_id_type>(get_group_id()[0]);
+ }
+
+ /// Equivalent to return get_local_id()[0].
+ linear_id_type get_local_linear_id() const noexcept {
+ return static_cast<linear_id_type>(get_local_id()[0]);
+ }
+
+ /// Equivalent to return get_group_range()[0].
+ linear_id_type get_group_linear_range() const noexcept {
+ return static_cast<linear_id_type>(get_group_range()[0]);
+ }
+
+ /// Equivalent to return get_local_range()[0].
+ linear_id_type get_local_linear_range() const noexcept {
+ return static_cast<linear_id_type>(get_local_range()[0]);
+ }
+
+ /// \return true for exactly one work-item in the sub-group, if the calling
+ /// work-item is the leader of the sub-group, and false for all other
+ /// work-items in the sub-group.
+ bool leader() const noexcept { return get_local_linear_id() == 0; }
+
+protected:
+ sub_group() = default;
+
+ template <int dimensions> friend class sycl::nd_item;
+};
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_SUB_GROUP_HPP
diff --git a/libsycl/include/sycl/__spirv/spirv_vars.hpp b/libsycl/include/sycl/__spirv/spirv_vars.hpp
index 450f581d9506d..854ae64b0f665 100644
--- a/libsycl/include/sycl/__spirv/spirv_vars.hpp
+++ b/libsycl/include/sycl/__spirv/spirv_vars.hpp
@@ -23,6 +23,16 @@
__attribute__((const)) size_t __spirv_BuiltInGlobalInvocationId(int);
__attribute__((const)) size_t __spirv_BuiltInGlobalSize(int);
__attribute__((const)) size_t __spirv_BuiltInGlobalOffset(int);
+__attribute__((const)) size_t __spirv_BuiltInWorkgroupId(int);
+__attribute__((const)) size_t __spirv_BuiltInLocalInvocationId(int);
+__attribute__((const)) size_t __spirv_BuiltInWorkgroupSize(int);
+__attribute__((const)) size_t __spirv_BuiltInNumWorkgroups(int);
+
+__attribute__((const)) uint32_t __spirv_BuiltInSubgroupSize();
+__attribute__((const)) uint32_t __spirv_BuiltInSubgroupMaxSize();
+__attribute__((const)) uint32_t __spirv_BuiltInNumSubgroups();
+__attribute__((const)) uint32_t __spirv_BuiltInSubgroupId();
+__attribute__((const)) uint32_t __spirv_BuiltInSubgroupLocalInvocationId();
namespace __spirv {
@@ -57,6 +67,10 @@ namespace __spirv {
__SPIRV_DEFINE_INIT_AND_GET_HELPERS(BuiltInGlobalSize);
__SPIRV_DEFINE_INIT_AND_GET_HELPERS(BuiltInGlobalInvocationId)
__SPIRV_DEFINE_INIT_AND_GET_HELPERS(BuiltInGlobalOffset)
+__SPIRV_DEFINE_INIT_AND_GET_HELPERS(BuiltInWorkgroupId)
+__SPIRV_DEFINE_INIT_AND_GET_HELPERS(BuiltInLocalInvocationId)
+__SPIRV_DEFINE_INIT_AND_GET_HELPERS(BuiltInWorkgroupSize)
+__SPIRV_DEFINE_INIT_AND_GET_HELPERS(BuiltInNumWorkgroups)
#undef __SPIRV_DEFINE_INIT_AND_GET_HELPERS
diff --git a/libsycl/include/sycl/sycl.hpp b/libsycl/include/sycl/sycl.hpp
index 7e81d952bd41c..597c3377be27d 100644
--- a/libsycl/include/sycl/sycl.hpp
+++ b/libsycl/include/sycl/sycl.hpp
@@ -19,9 +19,14 @@
#include <sycl/__impl/device_selector.hpp>
#include <sycl/__impl/event.hpp>
#include <sycl/__impl/exception.hpp>
+#include <sycl/__impl/group.hpp>
#include <sycl/__impl/index_space_classes.hpp>
+#include <sycl/__impl/memory_enums.hpp>
+#include <sycl/__impl/nd_item.hpp>
+#include <sycl/__impl/nd_range.hpp>
#include <sycl/__impl/platform.hpp>
#include <sycl/__impl/queue.hpp>
+#include <sycl/__impl/sub_group.hpp>
#include <sycl/__impl/usm_functions.hpp>
#endif // _LIBSYCL_SYCL_HPP
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index beba3082f06e2..ff213d4df66a5 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -38,6 +38,18 @@ static void setKernelLaunchArgs(const detail::UnifiedRangeView &Range,
}
}
+ // We have the following mapping between dimensions with SPIR-V builtins:
+ // 1D: id[0] -> x
+ // 2D: id[0] -> y, id[1] -> x
+ // 3D: id[0] -> z, id[1] -> y, id[2] -> x
+ // So in order to ensure the correctness we update all the kernel
+ // parameters accordingly.
+ if (Range.MDims > 1) {
+ // TODO: Offset is not supported in liboffload so just ignore it for now.
+ std::swap(GlobalSize[0], GlobalSize[Range.MDims - 1]);
+ std::swap(GroupSize[0], GroupSize[Range.MDims - 1]);
+ }
+
ArgsToSet.Dimensions = Range.MDims;
ArgsToSet.NumGroups.x = GlobalSize[0] / GroupSize[0];
ArgsToSet.NumGroups.y = GlobalSize[1] / GroupSize[1];
diff --git a/libsycl/test/basic/group.cpp b/libsycl/test/basic/group.cpp
new file mode 100644
index 0000000000000..cfc3870b5cf99
--- /dev/null
+++ b/libsycl/test/basic/group.cpp
@@ -0,0 +1,78 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <cstddef>
+
+template <typename T> void initialize(T *Ptr, size_t Count, T Value) {
+ for (size_t I = 0; I < Count; ++I)
+ Ptr[I] = Value;
+}
+
+int main() {
+ sycl::queue Q;
+
+ constexpr int Dims = 3;
+ const sycl::range<Dims> LocalRange{2, 3, 1};
+ const sycl::range<Dims> GroupRange{1, 2, 3};
+ const sycl::range<Dims> GlobalRange = LocalRange * GroupRange;
+ const size_t DataLen = GlobalRange.size();
+
+ size_t *GroupRangeData = sycl::malloc_shared<size_t>(DataLen * Dims, Q);
+ size_t *GroupLinearIdData = sycl::malloc_shared<size_t>(DataLen * Dims, Q);
+ initialize(GroupRangeData, DataLen * Dims, static_cast<size_t>(0));
+ initialize(GroupLinearIdData, DataLen * Dims, static_cast<size_t>(0));
+
+ Q.parallel_for<class group_get_group_range_regression>(
+ sycl::nd_range<3>{GlobalRange, LocalRange}, [=](sycl::nd_item<3> It) {
+ const size_t Off = It.get_global_linear_id() * Dims;
+ const auto GR = It.get_group().get_group_range();
+ GroupRangeData[Off + 0] = GR[0];
+ GroupRangeData[Off + 1] = GR[1];
+ GroupRangeData[Off + 2] = GR[2];
+ });
+
+ Q.parallel_for<class group_get_group_linear_id_regression>(
+ sycl::nd_range<3>{GlobalRange, LocalRange}, [=](sycl::nd_item<3> It) {
+ const size_t Off = It.get_global_linear_id() * Dims;
+ const size_t LI = It.get_group().get_group_linear_id();
+ GroupLinearIdData[Off + 0] = LI;
+ GroupLinearIdData[Off + 1] = LI;
+ GroupLinearIdData[Off + 2] = LI;
+ });
+
+ Q.wait();
+
+ const size_t SizeZ = GlobalRange.get(0);
+ const size_t SizeY = GlobalRange.get(1);
+ const size_t SizeX = GlobalRange.get(2);
+
+ bool Fail = false;
+ for (size_t Z = 0; Z < SizeZ; ++Z) {
+ for (size_t Y = 0; Y < SizeY; ++Y) {
+ for (size_t X = 0; X < SizeX; ++X) {
+ const size_t Ind = Z * SizeX * SizeY + Y * SizeX + X;
+
+ const size_t Off = Ind * Dims;
+ Fail |= GroupRangeData[Off + 0] != GroupRange.get(0);
+ Fail |= GroupRangeData[Off + 1] != GroupRange.get(1);
+ Fail |= GroupRangeData[Off + 2] != GroupRange.get(2);
+
+ const sycl::id<3> GlobalId{Z, Y, X};
+ const sycl::id<3> GroupId = GlobalId / LocalRange;
+ const size_t GoldLinearId =
+ GroupId.get(0) * GroupRange.get(1) * GroupRange.get(2) +
+ GroupId.get(1) * GroupRange.get(2) + GroupId.get(2);
+ Fail |= GroupLinearIdData[Off + 0] != GoldLinearId;
+ Fail |= GroupLinearIdData[Off + 1] != GoldLinearId;
+ Fail |= GroupLinearIdData[Off + 2] != GoldLinearId;
+ }
+ }
+ }
+
+ sycl::free(GroupRangeData, Q);
+ sycl::free(GroupLinearIdData, Q);
+ return Fail;
+}
diff --git a/libsycl/test/basic/group_local_id.cpp b/libsycl/test/basic/group_local_id.cpp
new file mode 100644
index 0000000000000..499a45f55a91a
--- /dev/null
+++ b/libsycl/test/basic/group_local_id.cpp
@@ -0,0 +1,47 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <cassert>
+
+template <int Dims> class group_local_id_kernel;
+
+template <int Dims>
+bool runGroupLocalIdCase(sycl::queue &Q, sycl::nd_range<Dims> ExecRange,
+ size_t Count) {
+ int *Out = sycl::malloc_shared<int>(Count, Q);
+ for (size_t I = 0; I < Count; ++I)
+ Out[I] = -1;
+
+ Q.parallel_for<group_local_id_kernel<Dims>>(
+ ExecRange, [=](sycl::nd_item<Dims> Item) {
+ Out[Item.get_global_linear_id()] =
+ (Item.get_local_id() == Item.get_group().get_local_id());
+ });
+
+ Q.wait();
+
+ bool Match = true;
+ for (size_t I = 0; I < Count; ++I)
+ Match &= (Out[I] == 1);
+
+ sycl::free(Out, Q);
+ return !Match;
+}
+
+int main() {
+ sycl::queue Q;
+ constexpr size_t N = 8;
+
+ bool Failure = runGroupLocalIdCase<1>(
+ Q, sycl::nd_range<1>{sycl::range<1>{N}, sycl::range<1>{N}}, N);
+ Failure |= runGroupLocalIdCase<2>(
+ Q, sycl::nd_range<2>{sycl::range<2>{N, N}, sycl::range<2>{N, N}}, N * N);
+ Failure |= runGroupLocalIdCase<3>(
+ Q, sycl::nd_range<3>{sycl::range<3>{N, N, N}, sycl::range<3>{N, N, N}},
+ N * N * N);
+
+ return Failure;
+}
diff --git a/libsycl/test/basic/linear_sub_group.cpp b/libsycl/test/basic/linear_sub_group.cpp
new file mode 100644
index 0000000000000..4f7e88568cc36
--- /dev/null
+++ b/libsycl/test/basic/linear_sub_group.cpp
@@ -0,0 +1,37 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <cassert>
+
+int main() {
+ sycl::queue Q;
+
+ constexpr uint32_t Outer = 2;
+ constexpr uint32_t Inner = 8;
+ constexpr uint32_t Size = Outer * Inner;
+
+ int *Output = sycl::malloc_shared<int>(Size, Q);
+ for (uint32_t I = 0; I < Size; ++I)
+ Output[I] = -1;
+
+ Q.parallel_for<class linear_sub_group>(
+ sycl::nd_range<2>(sycl::range<2>(Outer, Inner),
+ sycl::range<2>(Outer, Inner)),
+ [=](sycl::nd_item<2> It) {
+ sycl::sub_group SG = It.get_sub_group();
+ Output[It.get_global_linear_id()] =
+ SG.get_group_linear_id() * SG.get_local_linear_range() +
+ SG.get_local_linear_id();
+ });
+
+ Q.wait();
+
+ for (uint32_t I = 0; I < Size; ++I)
+ assert(Output[I] == static_cast<int>(I));
+
+ sycl::free(Output, Q);
+ return 0;
+}
diff --git a/libsycl/test/basic/nd_range.cpp b/libsycl/test/basic/nd_range.cpp
new file mode 100644
index 0000000000000..358d5b1c4c060
--- /dev/null
+++ b/libsycl/test/basic/nd_range.cpp
@@ -0,0 +1,54 @@
+// RUN: %clangxx -fsycl %s -o %t.out -Wno-error=deprecated-declarations
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <iostream>
+
+int main() {
+ sycl::nd_range<1> one_dim_nd_range_offset({4}, {2}, {1});
+ assert(one_dim_nd_range_offset.get_global_range() == sycl::range<1>(4));
+ assert(one_dim_nd_range_offset.get_local_range() == sycl::range<1>(2));
+ assert(one_dim_nd_range_offset.get_group_range() == sycl::range<1>(2));
+ assert(one_dim_nd_range_offset.get_offset() == sycl::id<1>(1));
+ std::cout << "one_dim_nd_range_offset passed " << std::endl;
+
+ sycl::nd_range<2> two_dim_nd_range_offset({8, 16}, {4, 8}, {1, 1});
+ assert(two_dim_nd_range_offset.get_global_range() == sycl::range<2>(8, 16));
+ assert(two_dim_nd_range_offset.get_local_range() == sycl::range<2>(4, 8));
+ assert(two_dim_nd_range_offset.get_group_range() == sycl::range<2>(2, 2));
+ assert(two_dim_nd_range_offset.get_offset() == sycl::id<2>(1, 1));
+ std::cout << "two_dim_nd_range_offset passed " << std::endl;
+
+ sycl::nd_range<3> three_dim_nd_range_offset({32, 64, 128}, {16, 32, 64},
+ {1, 1, 1});
+ assert(three_dim_nd_range_offset.get_global_range() ==
+ sycl::range<3>(32, 64, 128));
+ assert(three_dim_nd_range_offset.get_local_range() ==
+ sycl::range<3>(16, 32, 64));
+ assert(three_dim_nd_range_offset.get_group_range() ==
+ sycl::range<3>(2, 2, 2));
+ assert(three_dim_nd_range_offset.get_offset() == sycl::id<3>(1, 1, 1));
+ std::cout << "three_dim_nd_range_offset passed " << std::endl;
+
+ sycl::nd_range<1> one_dim_nd_range({4}, {2});
+ assert(one_dim_nd_range.get_global_range() == sycl::range<1>(4));
+ assert(one_dim_nd_range.get_local_range() == sycl::range<1>(2));
+ assert(one_dim_nd_range.get_group_range() == sycl::range<1>(2));
+ assert(one_dim_nd_range.get_offset() == sycl::id<1>(0));
+ std::cout << "one_dim_nd_range passed " << std::endl;
+
+ sycl::nd_range<2> two_dim_nd_range({8, 16}, {4, 8});
+ assert(two_dim_nd_range.get_global_range() == sycl::range<2>(8, 16));
+ assert(two_dim_nd_range.get_local_range() == sycl::range<2>(4, 8));
+ assert(two_dim_nd_range.get_group_range() == sycl::range<2>(2, 2));
+ assert(two_dim_nd_range.get_offset() == sycl::id<2>(0, 0));
+ std::cout << "two_dim_nd_range passed " << std::endl;
+
+ sycl::nd_range<3> three_dim_nd_range({32, 64, 128}, {16, 32, 64});
+ assert(three_dim_nd_range.get_global_range() == sycl::range<3>(32, 64, 128));
+ assert(three_dim_nd_range.get_local_range() == sycl::range<3>(16, 32, 64));
+ assert(three_dim_nd_range.get_group_range() == sycl::range<3>(2, 2, 2));
+ assert(three_dim_nd_range.get_offset() == sycl::id<3>(0, 0, 0));
+ std::cout << "three_dim_nd_range passed " << std::endl;
+}
diff --git a/libsycl/test/basic/queue_parallel_for_generic.cpp b/libsycl/test/basic/queue_parallel_for_generic.cpp
index 70a191e6ab220..03943f7ed141f 100644
--- a/libsycl/test/basic/queue_parallel_for_generic.cpp
+++ b/libsycl/test/basic/queue_parallel_for_generic.cpp
@@ -47,14 +47,32 @@ int main() {
A[i.get_linear_id()]++;
});
- // TODO: add kernel with offset and kernel with nd_range once they
- // are implemented.
+ sycl::nd_range<1> NDR(sycl::range<1>{N}, sycl::range<1>{2});
+ Q.parallel_for<class NdRange1D>(NDR, [=](auto nd_i) {
+ static_assert(std::is_same<decltype(nd_i), sycl::nd_item<1>>::value,
+ "lambda arg type is unexpected");
+ A[nd_i.get_global_id()]++;
+ });
+
+ sycl::nd_range<2> NDR2D(sycl::range<2>{4, 2}, sycl::range<2>{2, 1});
+ Q.parallel_for<class NdRange2D>(NDR2D, [=](auto nd_i) {
+ static_assert(std::is_same<decltype(nd_i), sycl::nd_item<2>>::value,
+ "lambda arg type is unexpected");
+ A[nd_i.get_global_linear_id()]++;
+ });
+
+ sycl::nd_range<3> NDR3D(sycl::range<3>{2, 2, 2}, sycl::range<3>{1, 2, 2});
+ Q.parallel_for<class NdRange3D>(NDR3D, [=](auto nd_i) {
+ static_assert(std::is_same<decltype(nd_i), sycl::nd_item<3>>::value,
+ "lambda arg type is unexpected");
+ A[nd_i.get_global_linear_id()]++;
+ });
Q.wait();
bool Fail{};
for (int i = 0; i < N; i++) {
- Fail |= !(A[i] == 5);
+ Fail |= !(A[i] == 8);
}
sycl::free(A, Ctx);
return Fail;
diff --git a/libsycl/test/basic/sub_group_by_value_semantics.cpp b/libsycl/test/basic/sub_group_by_value_semantics.cpp
new file mode 100644
index 0000000000000..9db3e3044fd47
--- /dev/null
+++ b/libsycl/test/basic/sub_group_by_value_semantics.cpp
@@ -0,0 +1,39 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+int main() {
+ sycl::queue Q;
+ bool *Result = sycl::malloc_shared<bool>(1, Q);
+ Result[0] = true;
+
+ Q.parallel_for<class sub_group_by_value_semantics>(
+ sycl::nd_range<3>({1, 1, 1}, {1, 1, 1}), [=](sycl::nd_item<3> Item) {
+ sycl::sub_group A = Item.get_sub_group();
+
+ // Check reflexivity.
+ Result[0] &= (A == A);
+ Result[0] &= !(A != A);
+
+ // Check symmetry.
+ auto Copied = A;
+ auto &B = Copied;
+ Result[0] &= (A == B);
+ Result[0] &= (B == A);
+ Result[0] &= !(A != B);
+ Result[0] &= !(B != A);
+
+ // Check transitivity.
+ auto CopiedTwice = Copied;
+ const auto &C = CopiedTwice;
+ Result[0] &= (C == A);
+ });
+
+ Q.wait();
+
+ bool Fail = !Result[0];
+ sycl::free(Result, Q);
+ return Fail;
+}
diff --git a/libsycl/test/basic/sub_group_common.cpp b/libsycl/test/basic/sub_group_common.cpp
new file mode 100644
index 0000000000000..11ab99248cc55
--- /dev/null
+++ b/libsycl/test/basic/sub_group_common.cpp
@@ -0,0 +1,80 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+struct Data {
+ unsigned int LocalId;
+ unsigned int LocalRange;
+ unsigned int MaxLocalRange;
+ unsigned int GroupId;
+ unsigned int GroupRange;
+};
+
+bool check(sycl::queue &Q, unsigned int G, unsigned int L) {
+ Data *SyclData = sycl::malloc_shared<Data>(G, Q);
+ size_t *SgSize = sycl::malloc_shared<size_t>(1, Q);
+
+ for (unsigned int I = 0; I < G; ++I)
+ SyclData[I] = {0, 0, 0, 0, 0};
+ SgSize[0] = 0;
+
+ Q.parallel_for<class sycl_subgr_common>(
+ sycl::nd_range<1>(sycl::range<1>(G), sycl::range<1>(L)),
+ [=](sycl::nd_item<1> NdItem) {
+ sycl::sub_group SG = NdItem.get_sub_group();
+ const unsigned int Index =
+ static_cast<unsigned int>(NdItem.get_global_id(0));
+ SyclData[Index].LocalId =
+ static_cast<unsigned int>(SG.get_local_id()[0]);
+ SyclData[Index].LocalRange =
+ static_cast<unsigned int>(SG.get_local_range()[0]);
+ SyclData[Index].MaxLocalRange =
+ static_cast<unsigned int>(SG.get_max_local_range()[0]);
+ SyclData[Index].GroupId =
+ static_cast<unsigned int>(SG.get_group_id()[0]);
+ SyclData[Index].GroupRange =
+ static_cast<unsigned int>(SG.get_group_range()[0]);
+ if (Index == 0)
+ SgSize[0] = SG.get_max_local_range()[0];
+ });
+
+ Q.wait();
+
+ bool Fail = false;
+ const unsigned int SGSize = static_cast<unsigned int>(SgSize[0]);
+ if (SGSize == 0) {
+ sycl::free(SyclData, Q);
+ sycl::free(SgSize, Q);
+ return true;
+ }
+
+ const unsigned int NumSg = L / SGSize + ((L % SGSize) ? 1U : 0U);
+ for (unsigned int J = 0; J < G; ++J) {
+ const unsigned int GroupId = (J % L) / SGSize;
+ const unsigned int LocalRange =
+ (GroupId + 1 == NumSg) ? (L - GroupId * SGSize) : SGSize;
+ Fail |= (SyclData[J].LocalId != ((J % L) % SGSize));
+ Fail |= (SyclData[J].LocalRange != LocalRange);
+ Fail |= (SyclData[J].MaxLocalRange != SyclData[0].MaxLocalRange);
+ Fail |= (SyclData[J].GroupId != GroupId);
+ Fail |= (SyclData[J].GroupRange != NumSg);
+ }
+
+ sycl::free(SyclData, Q);
+ sycl::free(SgSize, Q);
+ return Fail;
+}
+
+int main() {
+ sycl::queue Q;
+ bool Fail = false;
+
+ Fail |= check(Q, 240, 80);
+ Fail |= check(Q, 8, 4);
+ Fail |= check(Q, 24, 12);
+ Fail |= check(Q, 1024, 256);
+
+ return Fail;
+}
>From 7c472743374fefbda3d2a5ea8be6378aee2a1201 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Tue, 30 Jun 2026 05:29:34 -0700
Subject: [PATCH 2/5] fix comments
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/include/sycl/__impl/detail/kernel_arg_helpers.hpp | 6 ------
libsycl/include/sycl/__impl/memory_enums.hpp | 2 ++
libsycl/include/sycl/__impl/nd_item.hpp | 2 +-
libsycl/test/basic/nd_range.cpp | 1 +
4 files changed, 4 insertions(+), 7 deletions(-)
diff --git a/libsycl/include/sycl/__impl/detail/kernel_arg_helpers.hpp b/libsycl/include/sycl/__impl/detail/kernel_arg_helpers.hpp
index 4967f8657c28e..2fc7a82d078df 100644
--- a/libsycl/include/sycl/__impl/detail/kernel_arg_helpers.hpp
+++ b/libsycl/include/sycl/__impl/detail/kernel_arg_helpers.hpp
@@ -127,12 +127,6 @@ class Builder {
return __spirv::initBuiltInGlobalInvocationId<Dims, id<Dims>>();
}
- /// \return the work group currently being operated on by the device.
- template <int Dims> static const group<Dims> getElement(group<Dims> *) {
- static_assert(isValidDimensions<Dims>, "invalid dimensions");
- return group<Dims>();
- }
-
/// \return the nd_item currently being operated on by the device.
template <int Dims> static const nd_item<Dims> getElement(nd_item<Dims> *) {
static_assert(isValidDimensions<Dims>, "invalid dimensions");
diff --git a/libsycl/include/sycl/__impl/memory_enums.hpp b/libsycl/include/sycl/__impl/memory_enums.hpp
index 37fbb1111465d..c6fecf4c5557c 100644
--- a/libsycl/include/sycl/__impl/memory_enums.hpp
+++ b/libsycl/include/sycl/__impl/memory_enums.hpp
@@ -16,6 +16,8 @@
#include <sycl/__impl/detail/config.hpp>
+#include <cstdint>
+
_LIBSYCL_BEGIN_NAMESPACE_SYCL
enum class memory_scope : std::uint32_t {
diff --git a/libsycl/include/sycl/__impl/nd_item.hpp b/libsycl/include/sycl/__impl/nd_item.hpp
index 632269349c925..f9ce98a67deb0 100644
--- a/libsycl/include/sycl/__impl/nd_item.hpp
+++ b/libsycl/include/sycl/__impl/nd_item.hpp
@@ -43,7 +43,7 @@ template <int Dimensions = 1> class nd_item {
nd_item &operator=(const nd_item &rhs) = default;
nd_item &operator=(nd_item &&rhs) = default;
- friend bool operator==(const nd_item &lhs, const nd_item &rhs) {
+ friend bool operator==(const nd_item &, const nd_item &) {
// https://github.com/KhronosGroup/SYCL-Docs/issues/532
return true;
}
diff --git a/libsycl/test/basic/nd_range.cpp b/libsycl/test/basic/nd_range.cpp
index 358d5b1c4c060..7988e8f3509f8 100644
--- a/libsycl/test/basic/nd_range.cpp
+++ b/libsycl/test/basic/nd_range.cpp
@@ -3,6 +3,7 @@
#include <sycl/sycl.hpp>
+#include <cassert>
#include <iostream>
int main() {
>From d0609f952444f7b327a6d747ec0623ba8decc770 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Thu, 2 Jul 2026 09:50:53 -0700
Subject: [PATCH 3/5] add UT for dimensions swap
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
.../unittests/queue/sycl_kernel_launch.cpp | 96 +++++++++++++++++++
1 file changed, 96 insertions(+)
diff --git a/libsycl/unittests/queue/sycl_kernel_launch.cpp b/libsycl/unittests/queue/sycl_kernel_launch.cpp
index 653dc7306d623..9a60d4c8b919b 100644
--- a/libsycl/unittests/queue/sycl_kernel_launch.cpp
+++ b/libsycl/unittests/queue/sycl_kernel_launch.cpp
@@ -49,6 +49,7 @@ class ScopedKernelRegistration {
class sycl::detail::MockQueue : public sycl::queue {
public:
+ using sycl::queue::setKernelParameters;
using sycl::queue::sycl_kernel_launch;
};
@@ -96,3 +97,98 @@ TEST(Queue, KernelLaunch) {
EXPECT_CALL(Mock.get(), olSyncQueue(_)).Times(1);
Q.wait();
}
+
+// Captures the ol_kernel_launch_size_args_t passed to olLaunchKernel.
+// SetParams is called with the queue to invoke setKernelParameters.
+static ol_kernel_launch_size_args_t captureKernelLaunchArgs(
+ mock::MockWrapper &Mock,
+ std::function<void(sycl::detail::MockQueue &)> SetParams) {
+ ScopedKernelRegistration Reg{"DimSwapTestKernel"};
+ sycl::detail::MockQueue Q;
+ ol_kernel_launch_size_args_t Captured{};
+ EXPECT_CALL(Mock.get(), olLaunchKernel(_, _, _, _, _, 1, _, _))
+ .WillOnce([&Captured](ol_queue_handle_t, ol_device_handle_t,
+ ol_symbol_handle_t,
+ const ol_kernel_launch_size_args_t *Args,
+ const ol_kernel_launch_prop_t *, size_t, void **,
+ const size_t *) -> ol_result_t {
+ Captured = *Args;
+ return OL_SUCCESS;
+ });
+ SetParams(Q);
+ KernelData Data{};
+ Q.sycl_kernel_launch<class DimSwapTestKernel>("DimSwapTestKernel", Data);
+ return Captured;
+}
+
+struct DimSwapParam {
+ // Name used by the test runner to identify the case.
+ const char *Description;
+ // Calls setKernelParameters on Q with the appropriate range.
+ std::function<void(sycl::detail::MockQueue &)> SetParams;
+ // Expected fields of ol_kernel_launch_size_args_t after the swap.
+ uint32_t ExpDims;
+ uint32_t ExpNGx, ExpNGy, ExpNGz;
+ uint32_t ExpGSx, ExpGSy, ExpGSz;
+};
+
+class DimSwapTest : public ::testing::TestWithParam<DimSwapParam> {};
+
+// Verifies that setKernelLaunchArgs correctly maps SYCL range dimensions to
+// liboffload's x/y/z axes.
+TEST_P(DimSwapTest, CheckLaunchArgs) {
+ mock::MockWrapper Mock;
+ const auto &P = GetParam();
+ auto Args = captureKernelLaunchArgs(Mock, P.SetParams);
+
+ EXPECT_EQ(Args.Dimensions, P.ExpDims);
+ EXPECT_EQ(Args.NumGroups.x, P.ExpNGx);
+ EXPECT_EQ(Args.NumGroups.y, P.ExpNGy);
+ EXPECT_EQ(Args.NumGroups.z, P.ExpNGz);
+ EXPECT_EQ(Args.GroupSize.x, P.ExpGSx);
+ EXPECT_EQ(Args.GroupSize.y, P.ExpGSy);
+ EXPECT_EQ(Args.GroupSize.z, P.ExpGSz);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+ DimensionSwap, DimSwapTest,
+ ::testing::Values(
+ // 1D nd_range: no swap.
+ // global={8}, local={2} -> x=4 groups of 2
+ DimSwapParam{"1D_NdRange",
+ [](sycl::detail::MockQueue &Q) {
+ sycl::nd_range<1> NDR(sycl::range<1>{8},
+ sycl::range<1>{2});
+ Q.setKernelParameters({}, NDR);
+ },
+ /*Dims=*/1, /*NG=*/4, 1, 1, /*GS=*/2, 1, 1},
+ // 2D nd_range: swap [0]<->[1].
+ // global={4,6}, local={2,3} -> after swap: global={6,4}, local={3,2}
+ DimSwapParam{"2D_NdRange",
+ [](sycl::detail::MockQueue &Q) {
+ sycl::nd_range<2> NDR(sycl::range<2>{4, 6},
+ sycl::range<2>{2, 3});
+ Q.setKernelParameters({}, NDR);
+ },
+ /*Dims=*/2, /*NG=*/2, 2, 1, /*GS=*/3, 2, 1},
+ // 3D nd_range: swap [0]<->[2].
+ // global={2,4,6}, local={1,2,3} -> after swap: global={6,4,2},
+ // local={3,2,1}
+ DimSwapParam{"3D_NdRange",
+ [](sycl::detail::MockQueue &Q) {
+ sycl::nd_range<3> NDR(sycl::range<3>{2, 4, 6},
+ sycl::range<3>{1, 2, 3});
+ Q.setKernelParameters({}, NDR);
+ },
+ /*Dims=*/3, /*NG=*/2, 2, 2, /*GS=*/3, 2, 1},
+ // 2D range (no local): swap [0]<->[1], GroupSize stays {1,1,1}.
+ // global={4,6} -> after swap: global={6,4}
+ DimSwapParam{"2D_Range",
+ [](sycl::detail::MockQueue &Q) {
+ sycl::range<2> R(4, 6);
+ Q.setKernelParameters({}, R);
+ },
+ /*Dims=*/2, /*NG=*/6, 4, 1, /*GS=*/1, 1, 1}),
+ [](const ::testing::TestParamInfo<DimSwapParam> &Info) {
+ return Info.param.Description;
+ });
>From b9e8aae603c8a8a5e45df3be2f865606a06c7ad7 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Tue, 14 Jul 2026 02:26:43 -0700
Subject: [PATCH 4/5] fix test
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/unittests/queue/sycl_kernel_launch.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/libsycl/unittests/queue/sycl_kernel_launch.cpp b/libsycl/unittests/queue/sycl_kernel_launch.cpp
index 9a60d4c8b919b..b23f5d7a5d601 100644
--- a/libsycl/unittests/queue/sycl_kernel_launch.cpp
+++ b/libsycl/unittests/queue/sycl_kernel_launch.cpp
@@ -103,7 +103,8 @@ TEST(Queue, KernelLaunch) {
static ol_kernel_launch_size_args_t captureKernelLaunchArgs(
mock::MockWrapper &Mock,
std::function<void(sycl::detail::MockQueue &)> SetParams) {
- ScopedKernelRegistration Reg{"DimSwapTestKernel"};
+ // Keep registration alive across all parameterized runs.
+ static ScopedKernelRegistration Reg{"DimSwapTestKernel"};
sycl::detail::MockQueue Q;
ol_kernel_launch_size_args_t Captured{};
EXPECT_CALL(Mock.get(), olLaunchKernel(_, _, _, _, _, 1, _, _))
>From 6578544897a2166dadad83cf31ef5ee31d809338 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Thu, 16 Jul 2026 03:10:58 -0700
Subject: [PATCH 5/5] fix code-review
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/include/sycl/__impl/group.hpp | 4 ++--
libsycl/include/sycl/__impl/queue.hpp | 2 +-
2 files changed, 3 insertions(+), 3 deletions(-)
diff --git a/libsycl/include/sycl/__impl/group.hpp b/libsycl/include/sycl/__impl/group.hpp
index edadcbf4cbacb..529ddd54f1400 100644
--- a/libsycl/include/sycl/__impl/group.hpp
+++ b/libsycl/include/sycl/__impl/group.hpp
@@ -28,7 +28,7 @@ template <int> class nd_item;
// SYCL2020 4.9.1.7. group class.
/// The group class encapsulates all functionality required to represent a
-/// particular work-group within a parallel execution.
+/// particular work-group within a parallel execution.
template <int Dimensions = 1> class group {
public:
using id_type = id<Dimensions>;
@@ -139,7 +139,7 @@ template <int Dimensions = 1> class group {
}
}
- template <int dimensions> friend class sycl::nd_item;
+ friend class sycl::nd_item<Dimensions>;
};
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index a4b9b30cdd172..da38b40d80148 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -419,7 +419,7 @@ class _LIBSYCL_EXPORT queue {
static_assert(
std::is_convertible_v<sycl::nd_item<Dims>, LambdaArgType>,
"Kernel argument of a sycl::parallel_for with sycl::nd_range "
- "must be sycl::nd_item");
+ "must be sycl::nd_item or be convertible from sycl::nd_item");
} else {
static_assert(
std::is_convertible_v<sycl::item<Dims>, LambdaArgType> ||
More information about the llvm-commits
mailing list