[libsycl] add operators to sycl::range and sycl::id (#203572)
This PR was assisted by GH Copilot (tests extension). Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova@intel.com>
This commit is contained in:
Kseniya Tikhomirova
GitHub
parent
84c42fe09e
commit
0bb2b00bee
4 changed files with 481 additions and 35 deletions
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -27,36 +27,40 @@ namespace detail {
|
|||
|
||||
class Builder;
|
||||
|
||||
template <typename N, typename T>
|
||||
using IntegralType = std::enable_if_t<std::is_integral_v<N>, T>;
|
||||
|
||||
/// Helper class for dimensions data management.
|
||||
template <int Dimensions = 1> class RawArray {
|
||||
template <typename Derived, int Dimensions> class IndexSpaceBase {
|
||||
static_assert(Dimensions >= 1 && Dimensions <= 3,
|
||||
"RawArray can only be 1, 2, or 3 Dimensional.");
|
||||
"IndexSpaceBase can only be 1, 2, or 3 Dimensional.");
|
||||
|
||||
public:
|
||||
/// Constructs a one-dimensional instance and assigns the corresponding data
|
||||
/// to Dim0 value. Available only if Dimensions = 1.
|
||||
template <int N = Dimensions, std::enable_if_t<N == 1, bool> = true>
|
||||
RawArray(size_t Dim0 = 0) : MArray{Dim0} {}
|
||||
IndexSpaceBase(size_t Dim0 = 0) : MArray{Dim0} {}
|
||||
|
||||
/// Constructs a two-dimensional instance and assigns the corresponding data.
|
||||
/// Available only if Dimensions = 2.
|
||||
template <int N = Dimensions, std::enable_if_t<N == 2, bool> = true>
|
||||
RawArray(size_t Dim0, size_t Dim1) : MArray{Dim0, Dim1} {}
|
||||
IndexSpaceBase(size_t Dim0, size_t Dim1) : MArray{Dim0, Dim1} {}
|
||||
|
||||
/// Constructs a two-dimensional instance with the zero-initialized
|
||||
/// corresponding data. Available only if Dimensions = 2.
|
||||
template <int N = Dimensions, std::enable_if_t<N == 2, bool> = true>
|
||||
RawArray() : RawArray(0, 0) {}
|
||||
IndexSpaceBase() : IndexSpaceBase(0, 0) {}
|
||||
|
||||
/// Constructs a three-dimensional instance and assigns the corresponding
|
||||
/// data. Available only if Dimensions = 3.
|
||||
template <int N = Dimensions, std::enable_if_t<N == 3, bool> = true>
|
||||
RawArray(size_t Dim0, size_t Dim1, size_t Dim2) : MArray{Dim0, Dim1, Dim2} {}
|
||||
IndexSpaceBase(size_t Dim0, size_t Dim1, size_t Dim2)
|
||||
: MArray{Dim0, Dim1, Dim2} {}
|
||||
|
||||
/// Constructs a three-dimensional instance with the zero-initialized
|
||||
/// corresponding data. Available only if Dimensions = 3.
|
||||
template <int N = Dimensions, std::enable_if_t<N == 3, bool> = true>
|
||||
RawArray() : RawArray(0, 0, 0) {}
|
||||
IndexSpaceBase() : IndexSpaceBase(0, 0, 0) {}
|
||||
|
||||
/// Returns the value for the specified dimension.
|
||||
/// Results in undefined behavior if dimension is not in the range [0,
|
||||
|
|
@ -81,14 +85,16 @@ public:
|
|||
return MArray[Dimension];
|
||||
}
|
||||
|
||||
RawArray(const RawArray<Dimensions> &rhs) = default;
|
||||
RawArray(RawArray<Dimensions> &&rhs) = default;
|
||||
RawArray<Dimensions> &operator=(const RawArray<Dimensions> &rhs) = default;
|
||||
RawArray<Dimensions> &operator=(RawArray<Dimensions> &&rhs) = default;
|
||||
~RawArray() = default;
|
||||
IndexSpaceBase(const IndexSpaceBase<Derived, Dimensions> &rhs) = default;
|
||||
IndexSpaceBase(IndexSpaceBase<Derived, Dimensions> &&rhs) = default;
|
||||
IndexSpaceBase<Derived, Dimensions> &
|
||||
operator=(const IndexSpaceBase<Derived, Dimensions> &rhs) = default;
|
||||
IndexSpaceBase<Derived, Dimensions> &
|
||||
operator=(IndexSpaceBase<Derived, Dimensions> &&rhs) = default;
|
||||
~IndexSpaceBase() = default;
|
||||
|
||||
friend bool operator==(const RawArray<Dimensions> &lhs,
|
||||
const RawArray<Dimensions> &rhs) {
|
||||
friend bool operator==(const IndexSpaceBase<Derived, Dimensions> &lhs,
|
||||
const IndexSpaceBase<Derived, Dimensions> &rhs) {
|
||||
for (int i = 0; i < Dimensions; ++i) {
|
||||
if (lhs.MArray[i] != rhs.MArray[i]) {
|
||||
return false;
|
||||
|
|
@ -97,14 +103,127 @@ public:
|
|||
return true;
|
||||
}
|
||||
|
||||
friend bool operator!=(const RawArray<Dimensions> &lhs,
|
||||
const RawArray<Dimensions> &rhs) {
|
||||
friend bool operator!=(const IndexSpaceBase<Derived, Dimensions> &lhs,
|
||||
const IndexSpaceBase<Derived, Dimensions> &rhs) {
|
||||
return !(lhs == rhs);
|
||||
}
|
||||
|
||||
#define _LIBSYCL_GEN_OPT(op) \
|
||||
friend Derived operator op(const Derived &lhs, \
|
||||
const Derived &rhs) noexcept { \
|
||||
Derived result; \
|
||||
for (int i = 0; i < Dimensions; ++i) { \
|
||||
result.MArray[i] = lhs.MArray[i] op rhs.MArray[i]; \
|
||||
} \
|
||||
return result; \
|
||||
} \
|
||||
\
|
||||
template <typename T> \
|
||||
friend IntegralType<T, Derived> operator op(const Derived &lhs, \
|
||||
const T &rhs) noexcept { \
|
||||
Derived result; \
|
||||
for (int i = 0; i < Dimensions; ++i) { \
|
||||
result.MArray[i] = lhs.MArray[i] op rhs; \
|
||||
} \
|
||||
return result; \
|
||||
} \
|
||||
\
|
||||
template <typename T> \
|
||||
friend IntegralType<T, Derived> operator op(const T &lhs, \
|
||||
const Derived &rhs) noexcept { \
|
||||
Derived 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
|
||||
|
||||
#define _LIBSYCL_GEN_OPT(op) \
|
||||
friend Derived &operator op(Derived &lhs, const Derived &rhs) noexcept { \
|
||||
for (int i = 0; i < Dimensions; ++i) { \
|
||||
lhs.MArray[i] op rhs[i]; \
|
||||
} \
|
||||
return lhs; \
|
||||
} \
|
||||
template <typename T> \
|
||||
friend IntegralType<T, Derived> &operator op(Derived &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
|
||||
|
||||
#define _LIBSYCL_GEN_OPT(op) \
|
||||
friend Derived operator op(const Derived &rhs) noexcept { \
|
||||
Derived 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
|
||||
|
||||
#define _LIBSYCL_GEN_OPT(op) \
|
||||
friend Derived &operator op(Derived &rhs) noexcept { \
|
||||
for (int i = 0; i < Dimensions; ++i) { \
|
||||
op rhs.MArray[i]; \
|
||||
} \
|
||||
return rhs; \
|
||||
} \
|
||||
friend Derived operator op(Derived &lhs, int) noexcept { \
|
||||
Derived oldLhs(lhs); \
|
||||
for (int i = 0; i < Dimensions; ++i) { \
|
||||
op lhs.MArray[i]; \
|
||||
} \
|
||||
return oldLhs; \
|
||||
}
|
||||
|
||||
_LIBSYCL_GEN_OPT(++)
|
||||
_LIBSYCL_GEN_OPT(--)
|
||||
|
||||
#undef _LIBSYCL_GEN_OPT
|
||||
|
||||
protected:
|
||||
size_t MArray[Dimensions];
|
||||
};
|
||||
|
||||
} // namespace detail
|
||||
|
||||
/// SYCL 2020 4.9.1.1. range class.
|
||||
|
|
@ -112,10 +231,10 @@ protected:
|
|||
/// domain of either a single work-group in a parallel dispatch, or the overall
|
||||
/// Dimensions of the dispatch.
|
||||
template <int Dimensions = 1>
|
||||
class range : public detail::RawArray<Dimensions> {
|
||||
class range : public detail::IndexSpaceBase<range<Dimensions>, Dimensions> {
|
||||
static_assert(Dimensions >= 1 && Dimensions <= 3,
|
||||
"range can only be 1-, 2-, or 3-dimensional.");
|
||||
using Base = detail::RawArray<Dimensions>;
|
||||
using Base = detail::IndexSpaceBase<range<Dimensions>, Dimensions>;
|
||||
|
||||
public:
|
||||
static constexpr int dimensions = Dimensions;
|
||||
|
|
@ -142,7 +261,7 @@ public:
|
|||
: Base(dim0, dim1, dim2) {}
|
||||
|
||||
/*
|
||||
Declared and implemented in detail::RawArray:
|
||||
Declared and implemented in detail::IndexSpaceBase:
|
||||
std::size_t get(int dimension) const noexcept;
|
||||
std::size_t& operator[](int dimension) noexcept;
|
||||
std::size_t operator[](int dimension) const noexcept;
|
||||
|
|
@ -156,8 +275,6 @@ public:
|
|||
}
|
||||
return size;
|
||||
}
|
||||
|
||||
// TODO: operators to be added
|
||||
};
|
||||
|
||||
/// c++ deduction guides.
|
||||
|
|
@ -173,10 +290,11 @@ template <int Dimensions = 1, bool WithOffset = true> class item;
|
|||
/// id<int Dimensions> is a vector of Dimensions that is used to represent an id
|
||||
/// into a global or local range. It can be used as an index in an accessor of
|
||||
/// the same rank.
|
||||
template <int Dimensions = 1> class id : public detail::RawArray<Dimensions> {
|
||||
template <int Dimensions = 1>
|
||||
class id : public detail::IndexSpaceBase<id<Dimensions>, Dimensions> {
|
||||
static_assert(Dimensions >= 1 && Dimensions <= 3,
|
||||
"id can only be 1-, 2-, or 3-dimensional.");
|
||||
using Base = detail::RawArray<Dimensions>;
|
||||
using Base = detail::IndexSpaceBase<id<Dimensions>, Dimensions>;
|
||||
|
||||
// Helper class for conversion operator. Void type is not suitable. User
|
||||
// cannot even try to get address of the operator PrivateTag(). User
|
||||
|
|
@ -230,39 +348,67 @@ public:
|
|||
|
||||
/// Constructs an id from item.get_id().
|
||||
/// Only valid when the template parameter Dimensions is equal to 1.
|
||||
template <int N = Dimensions, std::enable_if_t<N == 1, bool> = true>
|
||||
id(const item<Dimensions> &item) noexcept : Base(item.get_id(0)) {}
|
||||
template <int N = Dimensions, bool WithOffset = true,
|
||||
std::enable_if_t<N == 1, bool> = true>
|
||||
id(const item<Dimensions, WithOffset> &item) noexcept
|
||||
: Base(item.get_id(0)) {}
|
||||
|
||||
/// Constructs an id from item.get_id().
|
||||
/// Only valid when the template parameter Dimensions is equal to 2.
|
||||
template <int N = Dimensions, std::enable_if_t<N == 2, bool> = true>
|
||||
id(const item<Dimensions> &item) noexcept
|
||||
template <int N = Dimensions, bool WithOffset = true,
|
||||
std::enable_if_t<N == 2, bool> = true>
|
||||
id(const item<Dimensions, WithOffset> &item) noexcept
|
||||
: Base(item.get_id(0), item.get_id(1)) {}
|
||||
|
||||
/// Constructs an id from item.get_id().
|
||||
/// Only valid when the template parameter Dimensions is equal to 3.
|
||||
template <int N = Dimensions, std::enable_if_t<N == 3, bool> = true>
|
||||
id(const item<Dimensions> &item) noexcept
|
||||
template <int N = Dimensions, bool WithOffset = true,
|
||||
std::enable_if_t<N == 3, bool> = true>
|
||||
id(const item<Dimensions, WithOffset> &item) noexcept
|
||||
: Base(item.get_id(0), item.get_id(1), item.get_id(2)) {}
|
||||
|
||||
/*
|
||||
Declared and implemented in detail::RawArray:
|
||||
Declared and implemented in detail::IndexSpaceBase:
|
||||
std::size_t get(int dimension) const noexcept;
|
||||
std::size_t& operator[](int dimension) noexcept;
|
||||
std::size_t operator[](int dimension) const noexcept;
|
||||
*/
|
||||
|
||||
// Template operator is not allowed because it disables further type
|
||||
// conversion. For example, the next code will not work in case of template
|
||||
// conversion:
|
||||
// int a = id<1>(value);
|
||||
// conversion. For example, the next code will not work in case of template
|
||||
// conversion: int a = id<1>(value);
|
||||
/// Returns the same value as get(0).
|
||||
/// Available only when: Dimensions == 1.
|
||||
operator EnableIfT<(Dimensions == 1), std::size_t>() const noexcept {
|
||||
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 over type conversion. In the case of non-template operators,
|
||||
// ambiguity appears: "id op size_t" may refer "size_t op size_t" or "id op
|
||||
// size_t". In the 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
|
||||
};
|
||||
|
||||
/// c++ deduction guides.
|
||||
|
|
|
|||
284
libsycl/test/basic/index_space_classes.cpp
Normal file
284
libsycl/test/basic/index_space_classes.cpp
Normal file
|
|
@ -0,0 +1,284 @@
|
|||
// REQUIRES: any-device
|
||||
// RUN: %clangxx -fsycl %s -o %t.out
|
||||
// RUN: %t.out
|
||||
//
|
||||
// Unified test for sycl::range and sycl::id covering all operators defined in
|
||||
// theirs base class, plus class-specific behaviour for each type.
|
||||
|
||||
#include <cassert>
|
||||
#include <sycl/sycl.hpp>
|
||||
#include <type_traits>
|
||||
|
||||
using sycl::detail::Builder;
|
||||
|
||||
// Helper to create values for any dimension from three components.
|
||||
// Only uses the required number of values based on dimension.
|
||||
template <template <int> class T, int Dim>
|
||||
T<Dim> makeValue(std::size_t a, std::size_t b = 0, std::size_t c = 0) {
|
||||
if constexpr (Dim == 1)
|
||||
return T<1>(a);
|
||||
else if constexpr (Dim == 2)
|
||||
return T<2>(a, b);
|
||||
else
|
||||
return T<3>(a, b, c);
|
||||
}
|
||||
|
||||
// Tests binary and compound operators for a specific dimension.
|
||||
template <template <int> class T, int Dim> void testBinaryOpsForDim() {
|
||||
const T<Dim> P = makeValue<T, Dim>(8, 9, 10);
|
||||
const T<Dim> Q = makeValue<T, Dim>(2, 3, 5);
|
||||
|
||||
// Binary operators: object op object.
|
||||
assert((P + Q) == (makeValue<T, Dim>(10, 12, 15)));
|
||||
assert((P - Q) == (makeValue<T, Dim>(6, 6, 5)));
|
||||
assert((P * Q) == (makeValue<T, Dim>(16, 27, 50)));
|
||||
assert((P / Q) == (makeValue<T, Dim>(4, 3, 2)));
|
||||
assert((P % Q) == (makeValue<T, Dim>(0, 0, 0)));
|
||||
assert((P << Q) == (makeValue<T, Dim>(32, 72, 320)));
|
||||
assert((P >> Q) == (makeValue<T, Dim>(2, 1, 0)));
|
||||
assert((P & Q) == (makeValue<T, Dim>(0, 1, 0)));
|
||||
assert((P | Q) == (makeValue<T, Dim>(10, 11, 15)));
|
||||
assert((P ^ Q) == (makeValue<T, Dim>(10, 10, 15)));
|
||||
assert((P && Q) == (makeValue<T, Dim>(1, 1, 1)));
|
||||
assert((P || Q) == (makeValue<T, Dim>(1, 1, 1)));
|
||||
assert((P < Q) == (makeValue<T, Dim>(0, 0, 0)));
|
||||
assert((P > Q) == (makeValue<T, Dim>(1, 1, 1)));
|
||||
assert((P <= Q) == (makeValue<T, Dim>(0, 0, 0)));
|
||||
assert((P >= Q) == (makeValue<T, Dim>(1, 1, 1)));
|
||||
|
||||
// Binary operators: object op scalar.
|
||||
assert((P + 1) == (makeValue<T, Dim>(9, 10, 11)));
|
||||
assert((P - 1) == (makeValue<T, Dim>(7, 8, 9)));
|
||||
assert((P * 2) == (makeValue<T, Dim>(16, 18, 20)));
|
||||
assert((P / 2) == (makeValue<T, Dim>(4, 4, 5)));
|
||||
assert((P % 4) == (makeValue<T, Dim>(0, 1, 2)));
|
||||
assert((P << 1) == (makeValue<T, Dim>(16, 18, 20)));
|
||||
assert((P >> 1) == (makeValue<T, Dim>(4, 4, 5)));
|
||||
assert((P & 6) == (makeValue<T, Dim>(0, 0, 2)));
|
||||
assert((P | 1) == (makeValue<T, Dim>(9, 9, 11)));
|
||||
assert((P ^ 3) == (makeValue<T, Dim>(11, 10, 9)));
|
||||
assert((P && 0) == (makeValue<T, Dim>(0, 0, 0)));
|
||||
assert((P || 0) == (makeValue<T, Dim>(1, 1, 1)));
|
||||
assert((P < 9) == (makeValue<T, Dim>(1, 0, 0)));
|
||||
assert((P > 9) == (makeValue<T, Dim>(0, 0, 1)));
|
||||
assert((P <= 9) == (makeValue<T, Dim>(1, 1, 0)));
|
||||
assert((P >= 9) == (makeValue<T, Dim>(0, 1, 1)));
|
||||
|
||||
// Binary operators: scalar op object.
|
||||
assert((1 + P) == (makeValue<T, Dim>(9, 10, 11)));
|
||||
assert((20 - Q) == (makeValue<T, Dim>(18, 17, 15)));
|
||||
assert((2 * P) == (makeValue<T, Dim>(16, 18, 20)));
|
||||
assert((20 / Q) == (makeValue<T, Dim>(10, 6, 4)));
|
||||
assert((33 % Q) == (makeValue<T, Dim>(1, 0, 3)));
|
||||
assert((1 << Q) == (makeValue<T, Dim>(4, 8, 32)));
|
||||
assert((256 >> Q) == (makeValue<T, Dim>(64, 32, 8)));
|
||||
assert((15 & Q) == (makeValue<T, Dim>(2, 3, 5)));
|
||||
assert((1 | Q) == (makeValue<T, Dim>(3, 3, 5)));
|
||||
assert((3 ^ Q) == (makeValue<T, Dim>(1, 0, 6)));
|
||||
assert((1 && Q) == (makeValue<T, Dim>(1, 1, 1)));
|
||||
assert((0 || Q) == (makeValue<T, Dim>(1, 1, 1)));
|
||||
assert((9 < P) == (makeValue<T, Dim>(0, 0, 1)));
|
||||
assert((9 > P) == (makeValue<T, Dim>(1, 0, 0)));
|
||||
assert((9 <= P) == (makeValue<T, Dim>(0, 1, 1)));
|
||||
assert((9 >= P) == (makeValue<T, Dim>(1, 1, 0)));
|
||||
|
||||
// Compound assignment: object op= object.
|
||||
T<Dim> E = makeValue<T, Dim>(8, 9, 10);
|
||||
E += Q;
|
||||
assert(E == (makeValue<T, Dim>(10, 12, 15)));
|
||||
E -= Q;
|
||||
assert(E == (makeValue<T, Dim>(8, 9, 10)));
|
||||
E *= Q;
|
||||
assert(E == (makeValue<T, Dim>(16, 27, 50)));
|
||||
E /= Q;
|
||||
assert(E == (makeValue<T, Dim>(8, 9, 10)));
|
||||
E %= Q;
|
||||
assert(E == (makeValue<T, Dim>(0, 0, 0)));
|
||||
|
||||
// Compound assignment: object op= scalar.
|
||||
T<Dim> F = makeValue<T, Dim>(8, 9, 10);
|
||||
F += 2;
|
||||
assert(F == (makeValue<T, Dim>(10, 11, 12)));
|
||||
F -= 2;
|
||||
assert(F == (makeValue<T, Dim>(8, 9, 10)));
|
||||
F *= 2;
|
||||
assert(F == (makeValue<T, Dim>(16, 18, 20)));
|
||||
F /= 2;
|
||||
assert(F == (makeValue<T, Dim>(8, 9, 10)));
|
||||
F %= 6;
|
||||
assert(F == (makeValue<T, Dim>(2, 3, 4)));
|
||||
F <<= 1;
|
||||
assert(F == (makeValue<T, Dim>(4, 6, 8)));
|
||||
F >>= 1;
|
||||
assert(F == (makeValue<T, Dim>(2, 3, 4)));
|
||||
F &= 6;
|
||||
assert(F == (makeValue<T, Dim>(2, 2, 4)));
|
||||
F |= 1;
|
||||
assert(F == (makeValue<T, Dim>(3, 3, 5)));
|
||||
F ^= 2;
|
||||
assert(F == (makeValue<T, Dim>(1, 1, 7)));
|
||||
}
|
||||
|
||||
template <template <int> class T> void testCommon() {
|
||||
// Default construction zero-initialises every dimension.
|
||||
T<1> Def1;
|
||||
T<2> Def2;
|
||||
T<3> Def3;
|
||||
assert(Def1[0] == 0);
|
||||
assert(Def2[0] == 0 && Def2[1] == 0);
|
||||
assert(Def3[0] == 0 && Def3[1] == 0 && Def3[2] == 0);
|
||||
|
||||
// Dimensional constructors and element access via get() and const [].
|
||||
T<1> A1(5);
|
||||
T<2> A2(3, 7);
|
||||
T<3> A3(8, 9, 10);
|
||||
assert(A1.get(0) == 5 && A1[0] == 5);
|
||||
assert(A2.get(0) == 3 && A2.get(1) == 7);
|
||||
assert(A3.get(0) == 8 && A3.get(1) == 9 && A3.get(2) == 10);
|
||||
|
||||
// dimensions static constexpr.
|
||||
static_assert(T<1>::dimensions == 1);
|
||||
static_assert(T<2>::dimensions == 2);
|
||||
static_assert(T<3>::dimensions == 3);
|
||||
|
||||
// Non-const operator[] write.
|
||||
T<2> Mutable(1, 2);
|
||||
Mutable[0] = 10;
|
||||
Mutable[1] = 20;
|
||||
assert(Mutable[0] == 10 && Mutable[1] == 20);
|
||||
|
||||
// Copy constructor and copy assignment.
|
||||
T<3> CopyCtor(A3);
|
||||
assert(CopyCtor == A3);
|
||||
T<3> CopyAssign(1, 1, 1);
|
||||
CopyAssign = A3;
|
||||
assert(CopyAssign == A3);
|
||||
|
||||
// Move constructor and move assignment.
|
||||
T<3> MoveSrc1(A3);
|
||||
T<3> MoveCtor(std::move(MoveSrc1));
|
||||
assert(MoveCtor == A3);
|
||||
T<3> MoveSrc2(A3);
|
||||
T<3> MoveAssign(1, 1, 1);
|
||||
MoveAssign = std::move(MoveSrc2);
|
||||
assert(MoveAssign == A3);
|
||||
|
||||
// Equality and inequality between same-type objects.
|
||||
assert((makeValue<T, 3>(8, 9, 10)) == (makeValue<T, 3>(8, 9, 10)));
|
||||
assert((makeValue<T, 3>(8, 9, 10)) != (makeValue<T, 3>(8, 9, 11)));
|
||||
|
||||
// Test binary and compound operators for all dimensions.
|
||||
testBinaryOpsForDim<T, 1>();
|
||||
testBinaryOpsForDim<T, 2>();
|
||||
testBinaryOpsForDim<T, 3>();
|
||||
|
||||
// Unary + and -.
|
||||
T<3> U(8, 9, 10);
|
||||
assert((+U) == (makeValue<T, 3>(8, 9, 10)));
|
||||
assert((-U) == (makeValue<T, 3>(static_cast<std::size_t>(-8),
|
||||
static_cast<std::size_t>(-9),
|
||||
static_cast<std::size_t>(-10))));
|
||||
|
||||
// Pre- and post-increment / decrement.
|
||||
T<3> Inc(8, 9, 10);
|
||||
assert((++Inc) == (makeValue<T, 3>(9, 10, 11)));
|
||||
assert((Inc++) == (makeValue<T, 3>(9, 10, 11)));
|
||||
assert(Inc == (makeValue<T, 3>(10, 11, 12)));
|
||||
assert((--Inc) == (makeValue<T, 3>(9, 10, 11)));
|
||||
assert((Inc--) == (makeValue<T, 3>(9, 10, 11)));
|
||||
assert(Inc == (makeValue<T, 3>(8, 9, 10)));
|
||||
}
|
||||
|
||||
void testRange() {
|
||||
testCommon<sycl::range>();
|
||||
|
||||
assert(sycl::range<1>(7).size() == 7);
|
||||
assert(sycl::range<2>(4, 5).size() == 20);
|
||||
assert(sycl::range<3>(2, 3, 4).size() == 24);
|
||||
assert(sycl::range<3>(0, 5, 3).size() == 0);
|
||||
|
||||
#ifdef __cpp_deduction_guides
|
||||
auto R1 = sycl::range(7);
|
||||
auto R2 = sycl::range(4, 5);
|
||||
auto R3 = sycl::range(2, 3, 4);
|
||||
static_assert(std::is_same_v<decltype(R1), sycl::range<1>>);
|
||||
static_assert(std::is_same_v<decltype(R2), sycl::range<2>>);
|
||||
static_assert(std::is_same_v<decltype(R3), sycl::range<3>>);
|
||||
assert(R1 == sycl::range<1>(7));
|
||||
assert(R2 == sycl::range<2>(4, 5));
|
||||
assert(R3 == sycl::range<3>(2, 3, 4));
|
||||
#endif
|
||||
}
|
||||
|
||||
void testId() {
|
||||
testCommon<sycl::id>();
|
||||
|
||||
// Construction from range.
|
||||
assert(sycl::id<1>(sycl::range<1>(2)) == sycl::id<1>(2));
|
||||
assert(sycl::id<2>(sycl::range<2>(4, 8)) == sycl::id<2>(4, 8));
|
||||
assert(sycl::id<3>(sycl::range<3>(16, 32, 64)) == sycl::id<3>(16, 32, 64));
|
||||
|
||||
// Construction from item (WithOffset=true) preserves the 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));
|
||||
|
||||
// Construction from item (WithOffset=false).
|
||||
sycl::item<1, false> Item1NoOffset = Builder::createItem<1, false>({4}, {2});
|
||||
sycl::item<2, false> Item2NoOffset =
|
||||
Builder::createItem<2, false>({8, 16}, {4, 8});
|
||||
sycl::item<3, false> Item3NoOffset =
|
||||
Builder::createItem<3, false>({32, 64, 128}, {16, 32, 64});
|
||||
assert(sycl::id<1>(Item1NoOffset) == sycl::id<1>(2));
|
||||
assert(sycl::id<2>(Item2NoOffset) == sycl::id<2>(4, 8));
|
||||
assert(sycl::id<3>(Item3NoOffset) == sycl::id<3>(16, 32, 64));
|
||||
|
||||
// Copy-initialization from item must also select id(item) constructor.
|
||||
sycl::id<1> IdFromItem1NoOffset = Item1NoOffset;
|
||||
sycl::id<2> IdFromItem2NoOffset = Item2NoOffset;
|
||||
sycl::id<3> IdFromItem3NoOffset = Item3NoOffset;
|
||||
assert(IdFromItem1NoOffset == sycl::id<1>(2));
|
||||
assert(IdFromItem2NoOffset == sycl::id<2>(4, 8));
|
||||
assert(IdFromItem3NoOffset == sycl::id<3>(16, 32, 64));
|
||||
|
||||
#ifdef __cpp_deduction_guides
|
||||
auto I1 = sycl::id(2);
|
||||
auto I2 = sycl::id(4, 8);
|
||||
auto I3 = sycl::id(16, 32, 64);
|
||||
static_assert(std::is_same_v<decltype(I1), sycl::id<1>>);
|
||||
static_assert(std::is_same_v<decltype(I2), sycl::id<2>>);
|
||||
static_assert(std::is_same_v<decltype(I3), sycl::id<3>>);
|
||||
assert(I1 == sycl::id<1>(2));
|
||||
assert(I2 == sycl::id<2>(4, 8));
|
||||
assert(I3 == sycl::id<3>(16, 32, 64));
|
||||
#endif
|
||||
|
||||
// Implicit conversion to size_t (id<1> only).
|
||||
sycl::id<1> OneDim(16);
|
||||
std::size_t S = OneDim;
|
||||
int I = OneDim;
|
||||
assert(S == 16 && I == 16);
|
||||
|
||||
// Scalar equality/inequality operators (id<1> only).
|
||||
const sycl::id<1> ConstId(10);
|
||||
assert(ConstId == 10);
|
||||
assert(10 == ConstId);
|
||||
assert(ConstId != 19);
|
||||
assert(19 != ConstId);
|
||||
assert(ConstId == 10u);
|
||||
assert(10u == ConstId);
|
||||
assert(ConstId != static_cast<short>(9));
|
||||
assert(static_cast<short>(9) != ConstId);
|
||||
}
|
||||
|
||||
int main() {
|
||||
testRange();
|
||||
testId();
|
||||
return 0;
|
||||
}
|
||||
16
libsycl/test/basic/index_space_classes_negative.cpp
Normal file
16
libsycl/test/basic/index_space_classes_negative.cpp
Normal file
|
|
@ -0,0 +1,16 @@
|
|||
// RUN: %clangxx -fsycl %s -fsyntax-only -Xclang -verify \
|
||||
// RUN: -Xclang -verify-ignore-unexpected=note
|
||||
|
||||
#include <sycl/sycl.hpp>
|
||||
|
||||
void testInvalidDimensionGatedAPI() {
|
||||
sycl::id<2> Id2(1, 2);
|
||||
std::size_t ScalarFromId2 = Id2; // expected-error {{no viable conversion}}
|
||||
(void)ScalarFromId2;
|
||||
|
||||
sycl::id<2> BadId2(7); // expected-error {{no matching constructor}}
|
||||
sycl::range<3> BadRange3(1, 2); // expected-error {{no matching constructor}}
|
||||
|
||||
(void)BadId2;
|
||||
(void)BadRange3;
|
||||
}
|
||||
Loading…
Reference in a new issue