[llvm] [libsycl] add operators to sycl::range and sycl::id (PR #203572)
Kseniya Tikhomirova via llvm-commits
llvm-commits at lists.llvm.org
Fri Jun 12 08:50:41 PDT 2026
https://github.com/KseniyaTikhomirova created https://github.com/llvm/llvm-project/pull/203572
None
>From b05c1e6a80eb64080e1d9fb58ad2b210829e5867 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Fri, 29 May 2026 01:46:07 -0700
Subject: [PATCH] [libsycl] add operators to sycl::range and sycl::id
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/docs/index.rst | 2 +-
.../sycl/__impl/index_space_classes.hpp | 274 +++++++++++++++++-
libsycl/test/basic/id.cpp | 71 +++++
libsycl/test/basic/range.cpp | 133 +++++++++
4 files changed, 477 insertions(+), 3 deletions(-)
create mode 100644 libsycl/test/basic/id.cpp
create mode 100644 libsycl/test/basic/range.cpp
diff --git a/libsycl/docs/index.rst b/libsycl/docs/index.rst
index 4e92a219163ca..a5d657f3e6404 100644
--- a/libsycl/docs/index.rst
+++ b/libsycl/docs/index.rst
@@ -126,7 +126,7 @@ TODO for added SYCL classes
* handle sub devices once they are implemented (blocked by liboffload support)
* ``event``: get_wait_list, get_info, get_profiling_info, wait_and_throw & default ctor are not implemented
-* ``range``, ``id`` - to add operators
+* ``range``, ``id`` - __SYCL_DISABLE_ID_TO_INT_CONV__ and __SYCL_ASSUME_ID_RANGE optimizations are not implemented
* general opens:
* define a way to report errors from object dtors
diff --git a/libsycl/include/sycl/__impl/index_space_classes.hpp b/libsycl/include/sycl/__impl/index_space_classes.hpp
index 823dcea062d15..cdef217d7341c 100644
--- a/libsycl/include/sycl/__impl/index_space_classes.hpp
+++ b/libsycl/include/sycl/__impl/index_space_classes.hpp
@@ -105,6 +105,9 @@ template <int Dimensions = 1> class RawArray {
protected:
size_t MArray[Dimensions];
};
+
+template <typename N, typename T>
+using IntegralType = std::enable_if_t<std::is_integral_v<N>, T>;
} // namespace detail
/// SYCL 2020 4.9.1.1. range class.
@@ -157,7 +160,124 @@ class range : public detail::RawArray<Dimensions> {
return size;
}
- // TODO: operators to be added
+ // OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=
+
+#define _LIBSYCL_GEN_OPT(op) \
+ friend range<Dimensions> operator op( \
+ const range<Dimensions> &lhs, const range<Dimensions> &rhs) noexcept { \
+ range<Dimensions> result(lhs); \
+ for (int i = 0; i < Dimensions; ++i) { \
+ result.MArray[i] = lhs.MArray[i] op rhs.MArray[i]; \
+ } \
+ return result; \
+ } \
+ \
+ template <typename T> \
+ friend detail::IntegralType<T, range<Dimensions>> operator op( \
+ const range<Dimensions> &lhs, const T &rhs) noexcept { \
+ range<Dimensions> result(lhs); \
+ for (int i = 0; i < Dimensions; ++i) { \
+ result.MArray[i] = lhs.MArray[i] op rhs; \
+ } \
+ return result; \
+ } \
+ \
+ template <typename T> \
+ friend detail::IntegralType<T, range<Dimensions>> operator op( \
+ const T &lhs, const range<Dimensions> &rhs) noexcept { \
+ range<Dimensions> result(rhs); \
+ for (int i = 0; i < Dimensions; ++i) { \
+ result.MArray[i] = lhs op rhs.MArray[i]; \
+ } \
+ return result; \
+ }
+
+ _LIBSYCL_GEN_OPT(+)
+ _LIBSYCL_GEN_OPT(-)
+ _LIBSYCL_GEN_OPT(*)
+ _LIBSYCL_GEN_OPT(/)
+ _LIBSYCL_GEN_OPT(%)
+ _LIBSYCL_GEN_OPT(<<)
+ _LIBSYCL_GEN_OPT(>>)
+ _LIBSYCL_GEN_OPT(&)
+ _LIBSYCL_GEN_OPT(|)
+ _LIBSYCL_GEN_OPT(^)
+ _LIBSYCL_GEN_OPT(&&)
+ _LIBSYCL_GEN_OPT(||)
+ _LIBSYCL_GEN_OPT(<)
+ _LIBSYCL_GEN_OPT(>)
+ _LIBSYCL_GEN_OPT(<=)
+ _LIBSYCL_GEN_OPT(>=)
+
+#undef _LIBSYCL_GEN_OPT
+
+// OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
+#define _LIBSYCL_GEN_OPT(op) \
+ friend range<Dimensions> &operator op( \
+ range<Dimensions> &lhs, const range<Dimensions> &rhs) noexcept { \
+ for (int i = 0; i < Dimensions; ++i) { \
+ lhs.MArray[i] op rhs[i]; \
+ } \
+ return lhs; \
+ } \
+ template <typename T> \
+ friend detail::IntegralType<T, range<Dimensions>> &operator op( \
+ range<Dimensions> &lhs, const T &rhs) noexcept { \
+ for (int i = 0; i < Dimensions; ++i) { \
+ lhs.MArray[i] op rhs; \
+ } \
+ return lhs; \
+ }
+
+ _LIBSYCL_GEN_OPT(+=)
+ _LIBSYCL_GEN_OPT(-=)
+ _LIBSYCL_GEN_OPT(*=)
+ _LIBSYCL_GEN_OPT(/=)
+ _LIBSYCL_GEN_OPT(%=)
+ _LIBSYCL_GEN_OPT(<<=)
+ _LIBSYCL_GEN_OPT(>>=)
+ _LIBSYCL_GEN_OPT(&=)
+ _LIBSYCL_GEN_OPT(|=)
+ _LIBSYCL_GEN_OPT(^=)
+
+#undef _LIBSYCL_GEN_OPT
+
+// OP is unary +, -
+#define _LIBSYCL_GEN_OPT(op) \
+ friend range<Dimensions> operator op( \
+ const range<Dimensions> &rhs) noexcept { \
+ range<Dimensions> result(rhs); \
+ for (int i = 0; i < Dimensions; ++i) { \
+ result.MArray[i] = (op rhs.MArray[i]); \
+ } \
+ return result; \
+ }
+
+ _LIBSYCL_GEN_OPT(+)
+ _LIBSYCL_GEN_OPT(-)
+
+#undef _LIBSYCL_GEN_OPT
+
+// OP is prefix ++, --
+#define _LIBSYCL_GEN_OPT(op) \
+ friend range<Dimensions> &operator op(range<Dimensions> &rhs) noexcept { \
+ for (int i = 0; i < Dimensions; ++i) { \
+ op rhs.MArray[i]; \
+ } \
+ return rhs; \
+ } \
+ friend range<Dimensions> operator op(range<Dimensions> &lhs, int) noexcept { \
+ range<Dimensions> old_lhs(lhs); \
+ for (int i = 0; i < Dimensions; ++i) { \
+ op lhs.MArray[i]; \
+ } \
+ return old_lhs; \
+ }
+
+ _LIBSYCL_GEN_OPT(++)
+ _LIBSYCL_GEN_OPT(--)
+
+#undef _LIBSYCL_GEN_OPT
};
/// c++ deduction guides.
@@ -262,7 +382,157 @@ template <int Dimensions = 1> class id : public detail::RawArray<Dimensions> {
return Base::get(0);
}
- // TODO: operators to be added
+// These operators are not a part of SYCL 2020 spec but are needed to avoid
+// ambiguity in case of implicit conversion id<1> vs size_t. Template
+// operators take precedence than type conversion. In the case of
+// non-template operators, ambiguity appears: "id op size_t" may refer
+// "size_t op size_t" and "id op size_t". In case of template operators it
+// will be "id op size_t".
+#define _LIBSYCL_GEN_OPT(op) \
+ template <typename T, int N = Dimensions, \
+ std::enable_if_t<N == 1, bool> = true> \
+ detail::IntegralType<T, bool> operator op(const T &rhs) const noexcept { \
+ if (this->MArray[0] != rhs) \
+ return false op true; \
+ return true op true; \
+ } \
+ template <typename T, int N = Dimensions, \
+ std::enable_if_t<N == 1, bool> = true> \
+ friend detail::IntegralType<T, bool> operator op( \
+ const T &lhs, const id<dimensions> &rhs) noexcept { \
+ if (lhs != rhs.MArray[0]) \
+ return false op true; \
+ return true op true; \
+ }
+
+ _LIBSYCL_GEN_OPT(==)
+ _LIBSYCL_GEN_OPT(!=)
+
+#undef _LIBSYCL_GEN_OPT
+
+// OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=
+#define _LIBSYCL_GEN_OPT(op) \
+ friend id<Dimensions> operator op(const id<Dimensions> &lhs, \
+ const id<Dimensions> &rhs) noexcept { \
+ id<Dimensions> result; \
+ for (int i = 0; i < Dimensions; ++i) { \
+ result.MArray[i] = lhs.MArray[i] op rhs.MArray[i]; \
+ } \
+ return result; \
+ } \
+ template <typename T> \
+ friend detail::IntegralType<T, id<Dimensions>> operator op( \
+ const id<Dimensions> &lhs, const T &rhs) noexcept { \
+ id<Dimensions> result; \
+ for (int i = 0; i < Dimensions; ++i) { \
+ result.MArray[i] = lhs.MArray[i] op rhs; \
+ } \
+ return result; \
+ } \
+ template <typename T> \
+ friend detail::IntegralType<T, id<Dimensions>> operator op( \
+ const T &lhs, const id<Dimensions> &rhs) noexcept { \
+ id<Dimensions> result; \
+ for (int i = 0; i < Dimensions; ++i) { \
+ result.MArray[i] = lhs op rhs.MArray[i]; \
+ } \
+ return result; \
+ }
+
+ _LIBSYCL_GEN_OPT(+)
+ _LIBSYCL_GEN_OPT(-)
+ _LIBSYCL_GEN_OPT(*)
+ _LIBSYCL_GEN_OPT(/)
+ _LIBSYCL_GEN_OPT(%)
+ _LIBSYCL_GEN_OPT(<<)
+ _LIBSYCL_GEN_OPT(>>)
+ _LIBSYCL_GEN_OPT(&)
+ _LIBSYCL_GEN_OPT(|)
+ _LIBSYCL_GEN_OPT(^)
+ _LIBSYCL_GEN_OPT(&&)
+ _LIBSYCL_GEN_OPT(||)
+ _LIBSYCL_GEN_OPT(<)
+ _LIBSYCL_GEN_OPT(>)
+ _LIBSYCL_GEN_OPT(<=)
+ _LIBSYCL_GEN_OPT(>=)
+
+#undef _LIBSYCL_GEN_OPT
+
+// OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
+#define _LIBSYCL_GEN_OPT(op) \
+ friend id<Dimensions> &operator op(id<Dimensions> &lhs, \
+ const id<Dimensions> &rhs) noexcept { \
+ for (int i = 0; i < Dimensions; ++i) { \
+ lhs.MArray[i] op rhs.MArray[i]; \
+ } \
+ return lhs; \
+ } \
+ template <typename T> \
+ friend detail::IntegralType<T, id<Dimensions>> &operator op( \
+ id<Dimensions> &lhs, const T &rhs) noexcept { \
+ for (int i = 0; i < Dimensions; ++i) { \
+ lhs.MArray[i] op rhs; \
+ } \
+ return lhs; \
+ }
+
+ _LIBSYCL_GEN_OPT(+=)
+ _LIBSYCL_GEN_OPT(-=)
+ _LIBSYCL_GEN_OPT(*=)
+ _LIBSYCL_GEN_OPT(/=)
+ _LIBSYCL_GEN_OPT(%=)
+ _LIBSYCL_GEN_OPT(<<=)
+ _LIBSYCL_GEN_OPT(>>=)
+ _LIBSYCL_GEN_OPT(&=)
+ _LIBSYCL_GEN_OPT(|=)
+ _LIBSYCL_GEN_OPT(^=)
+
+#undef _LIBSYCL_GEN_OPT
+
+// OP is unary +, -
+#define _LIBSYCL_GEN_OPT(op) \
+ friend id<Dimensions> operator op(const id<Dimensions> &rhs) noexcept { \
+ id<Dimensions> result; \
+ for (int i = 0; i < Dimensions; ++i) { \
+ result.MArray[i] = (op rhs.MArray[i]); \
+ } \
+ return result; \
+ }
+
+ _LIBSYCL_GEN_OPT(+)
+ _LIBSYCL_GEN_OPT(-)
+
+#undef _LIBSYCL_GEN_OPT
+
+// OP is prefix ++, --
+#define _LIBSYCL_GEN_OPT(op) \
+ friend id<Dimensions> &operator op(id<Dimensions> &rhs) noexcept { \
+ for (int i = 0; i < Dimensions; ++i) { \
+ op rhs.MArray[i]; \
+ } \
+ return rhs; \
+ }
+
+ _LIBSYCL_GEN_OPT(++)
+ _LIBSYCL_GEN_OPT(--)
+
+#undef _LIBSYCL_GEN_OPT
+
+// OP is postfix ++, --
+#define _LIBSYCL_GEN_OPT(op) \
+ friend id<Dimensions> operator op(id<Dimensions> &lhs, int) noexcept { \
+ id<Dimensions> old_lhs; \
+ for (int i = 0; i < Dimensions; ++i) { \
+ old_lhs.MArray[i] = lhs.MArray[i]; \
+ op lhs.MArray[i]; \
+ } \
+ return old_lhs; \
+ }
+
+ _LIBSYCL_GEN_OPT(++)
+ _LIBSYCL_GEN_OPT(--)
+
+#undef _LIBSYCL_GEN_OPT
};
/// c++ deduction guides.
diff --git a/libsycl/test/basic/id.cpp b/libsycl/test/basic/id.cpp
new file mode 100644
index 0000000000000..113814ee422ad
--- /dev/null
+++ b/libsycl/test/basic/id.cpp
@@ -0,0 +1,71 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <cassert>
+#include <sycl/sycl.hpp>
+
+using sycl::detail::Builder;
+
+int main() {
+ // Default construction initializes all dimensions to zero.
+ sycl::id<1> Zero1;
+ sycl::id<2> Zero2;
+ sycl::id<3> Zero3;
+ assert(Zero1 == sycl::id<1>(0));
+ assert(Zero2 == sycl::id<2>(0, 0));
+ assert(Zero3 == sycl::id<3>(0, 0, 0));
+
+ // Dimensional constructors and element access.
+ sycl::id<1> I1(64);
+ sycl::id<2> I2(128, 256);
+ sycl::id<3> I3(64, 1, 2);
+ assert(I1.get(0) == 64 && I1[0] == 64);
+ assert(I2.get(0) == 128 && I2.get(1) == 256);
+ assert(I3.get(0) == 64 && I3.get(1) == 1 && I3.get(2) == 2);
+
+ // Construction from range.
+ sycl::id<1> FromRange1(sycl::range<1>(2));
+ sycl::id<2> FromRange2(sycl::range<2>(4, 8));
+ sycl::id<3> FromRange3(sycl::range<3>(16, 32, 64));
+ assert(FromRange1 == sycl::id<1>(2));
+ assert(FromRange2 == sycl::id<2>(4, 8));
+ assert(FromRange3 == sycl::id<3>(16, 32, 64));
+
+ // Construction from item preserves computed id coordinates.
+ sycl::item<1, true> Item1 = Builder::createItem<1, true>({4}, {2}, {1});
+ sycl::item<2, true> Item2 =
+ Builder::createItem<2, true>({8, 16}, {4, 8}, {1, 1});
+ sycl::item<3, true> Item3 =
+ Builder::createItem<3, true>({32, 64, 128}, {16, 32, 64}, {1, 1, 1});
+ assert(sycl::id<1>(Item1) == sycl::id<1>(2));
+ assert(sycl::id<2>(Item2) == sycl::id<2>(4, 8));
+ assert(sycl::id<3>(Item3) == sycl::id<3>(16, 32, 64));
+
+ // Equality and inequality semantics.
+ assert(sycl::id<1>(10) == sycl::id<1>(10));
+ assert(sycl::id<2>(10, 15) != sycl::id<2>(10, 12));
+ assert(sycl::id<3>(1, 2, 3) == sycl::id<3>(1, 2, 3));
+ assert(sycl::id<1>(10) == 10);
+ assert(10 == sycl::id<1>(10));
+ assert(sycl::id<1>(10) != 19);
+
+ // 1D id implicit conversion to scalar.
+ sycl::id<1> OneDimCast(16);
+ std::size_t S = OneDimCast;
+ int I = OneDimCast;
+ assert(S == 16 && I == 16);
+
+ // Representative unary/increment behavior.
+ sycl::id<2> Unary(64, 1);
+ assert(+Unary == sycl::id<2>(64, 1));
+ assert(-Unary == sycl::id<2>(static_cast<std::size_t>(-64),
+ static_cast<std::size_t>(-1)));
+ assert(++Unary == sycl::id<2>(65, 2));
+ assert(Unary++ == sycl::id<2>(65, 2));
+ assert(Unary == sycl::id<2>(66, 3));
+ assert(--Unary == sycl::id<2>(65, 2));
+ assert(Unary-- == sycl::id<2>(65, 2));
+ assert(Unary == sycl::id<2>(64, 1));
+
+ return 0;
+}
diff --git a/libsycl/test/basic/range.cpp b/libsycl/test/basic/range.cpp
new file mode 100644
index 0000000000000..6a04f34cd7dcc
--- /dev/null
+++ b/libsycl/test/basic/range.cpp
@@ -0,0 +1,133 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <cassert>
+
+using namespace sycl;
+
+int main() {
+ sycl::range<1> OneDimRange(64);
+ sycl::range<2> TwoDimRange(64, 1);
+ sycl::range<3> ThreeDimRange(64, 1, 2);
+ assert(OneDimRange.size() == 64);
+ assert(OneDimRange.get(0) == 64);
+ assert(OneDimRange[0] == 64);
+ assert(TwoDimRange.size() == 64);
+ assert(TwoDimRange.get(0) == 64);
+ assert(TwoDimRange[0] == 64);
+ assert(TwoDimRange.get(1) == 1);
+ assert(TwoDimRange[1] == 1);
+ assert(ThreeDimRange.size() == 128);
+ assert(ThreeDimRange.get(0) == 64);
+ assert(ThreeDimRange[0] == 64);
+ assert(ThreeDimRange.get(1) == 1);
+ assert(ThreeDimRange[1] == 1);
+ assert(ThreeDimRange.get(2) == 2);
+ assert(ThreeDimRange[2] == 2);
+
+ sycl::range<3> Default3;
+ sycl::range<2> Default2;
+ sycl::range<1> Default1;
+
+ assert(Default3[0] == 0 && Default3[1] == 0 && Default3[2] == 0);
+ assert(Default2[0] == 0 && Default2[1] == 0);
+ assert(Default1[0] == 0);
+
+ const range<3> A(8, 9, 10);
+ const range<3> B(2, 3, 5);
+
+ assert((A + B) == range<3>(10, 12, 15));
+ assert((A - B) == range<3>(6, 6, 5));
+ assert((A * B) == range<3>(16, 27, 50));
+ assert((A / B) == range<3>(4, 3, 2));
+ assert((A % B) == range<3>(0, 0, 0));
+ assert((A << B) == range<3>(32, 72, 320));
+ assert((A >> B) == range<3>(2, 1, 0));
+ assert((A & B) == range<3>(0, 1, 0));
+ assert((A | B) == range<3>(10, 11, 15));
+ assert((A ^ B) == range<3>(10, 10, 15));
+ assert((A && B) == range<3>(1, 1, 1));
+ assert((A || B) == range<3>(1, 1, 1));
+ assert((A < B) == range<3>(0, 0, 0));
+ assert((A > B) == range<3>(1, 1, 1));
+ assert((A <= B) == range<3>(0, 0, 0));
+ assert((A >= B) == range<3>(1, 1, 1));
+
+ assert((A + 1) == range<3>(9, 10, 11));
+ assert((1 + A) == range<3>(9, 10, 11));
+ assert((A - 1) == range<3>(7, 8, 9));
+ assert((20 - B) == range<3>(18, 17, 15));
+ assert((A * 2) == range<3>(16, 18, 20));
+ assert((2 * A) == range<3>(16, 18, 20));
+ assert((A / 2) == range<3>(4, 4, 5));
+ assert((20 / B) == range<3>(10, 6, 4));
+ assert((A % 4) == range<3>(0, 1, 2));
+ assert((33 % B) == range<3>(1, 0, 3));
+ assert((A << 1) == range<3>(16, 18, 20));
+ assert((1 << B) == range<3>(4, 8, 32));
+ assert((A >> 1) == range<3>(4, 4, 5));
+ assert((256 >> B) == range<3>(64, 32, 8));
+ assert((A & 6) == range<3>(0, 0, 2));
+ assert((15 & B) == range<3>(2, 3, 5));
+ assert((A | 1) == range<3>(9, 9, 11));
+ assert((1 | B) == range<3>(3, 3, 5));
+ assert((A ^ 3) == range<3>(11, 10, 9));
+ assert((3 ^ B) == range<3>(1, 0, 6));
+ assert((A && 0) == range<3>(0, 0, 0));
+ assert((0 || B) == range<3>(1, 1, 1));
+ assert((A < 9) == range<3>(1, 0, 0));
+ assert((9 > A) == range<3>(1, 0, 0));
+ assert((A <= 9) == range<3>(1, 1, 0));
+ assert((9 >= A) == range<3>(1, 1, 0));
+
+ range<3> E = A;
+ E += B;
+ assert(E == range<3>(10, 12, 15));
+ E -= B;
+ assert(E == A);
+
+ E *= B;
+ assert(E == range<3>(16, 27, 50));
+ E /= B;
+ assert(E == A);
+ E %= B;
+ assert(E == range<3>(0, 0, 0));
+
+ range<3> F(8, 9, 10);
+ F += 2;
+ assert(F == range<3>(10, 11, 12));
+ F -= 2;
+ assert(F == range<3>(8, 9, 10));
+ F *= 2;
+ assert(F == range<3>(16, 18, 20));
+ F /= 2;
+ assert(F == range<3>(8, 9, 10));
+ F %= 6;
+ assert(F == range<3>(2, 3, 4));
+ F <<= 1;
+ assert(F == range<3>(4, 6, 8));
+ F >>= 1;
+ assert(F == range<3>(2, 3, 4));
+ F &= 6;
+ assert(F == range<3>(2, 2, 4));
+ F |= 1;
+ assert(F == range<3>(3, 3, 5));
+ F ^= 2;
+ assert(F == range<3>(1, 1, 7));
+
+ range<3> U(8, 9, 10);
+ assert(+U == range<3>(8, 9, 10));
+ assert(-U == range<3>(static_cast<std::size_t>(-8),
+ static_cast<std::size_t>(-9),
+ static_cast<std::size_t>(-10)));
+ assert(++U == range<3>(9, 10, 11));
+ assert(U++ == range<3>(9, 10, 11));
+ assert(U == range<3>(10, 11, 12));
+ assert(--U == range<3>(9, 10, 11));
+ assert(U-- == range<3>(9, 10, 11));
+ assert(U == range<3>(8, 9, 10));
+
+ return 0;
+}
More information about the llvm-commits
mailing list