diff --git a/sycl/include/CL/sycl/id.hpp b/sycl/include/CL/sycl/id.hpp index 1a993c281953e..f67855653391f 100644 --- a/sycl/include/CL/sycl/id.hpp +++ b/sycl/include/CL/sycl/id.hpp @@ -14,160 +14,261 @@ #include __SYCL_INLINE namespace cl { -namespace sycl { -template class range; -template class item; -template class id : public detail::array { -private: - using base = detail::array; - static_assert(dimensions >= 1 && dimensions <= 3, - "id can only be 1, 2, or 3 dimensional."); - template - using ParamTy = detail::enable_if_t<(N == val), T>; - -public: - id() = default; - - /* The following constructor is only available in the id struct - * specialization where: dimensions==1 */ - template id(ParamTy dim0) : base(dim0) {} - - template - id(ParamTy> &range_size) - : base(range_size.get(0)) {} - - template - id(ParamTy> &item) - : base(item.get_id(0)) {} - - /* The following constructor is only available in the id struct - * specialization where: dimensions==2 */ - template - id(ParamTy dim0, size_t dim1) : base(dim0, dim1) {} - - template - id(ParamTy> &range_size) - : base(range_size.get(0), range_size.get(1)) {} - - template - id(ParamTy> &item) - : base(item.get_id(0), item.get_id(1)) {} - - /* The following constructor is only available in the id struct - * specialization where: dimensions==3 */ - template - id(ParamTy dim0, size_t dim1, size_t dim2) - : base(dim0, dim1, dim2) {} - - template - id(ParamTy> &range_size) - : base(range_size.get(0), range_size.get(1), range_size.get(2)) {} - - template - id(ParamTy> &item) - : base(item.get_id(0), item.get_id(1), item.get_id(2)) {} - - explicit operator range() const { - range result( - detail::InitializedVal::template get<0>()); - for (int i = 0; i < dimensions; ++i) { - result[i] = this->get(i); + namespace sycl { + template class range; + template class item; + + template class id : public detail::array { + private: + using base = detail::array; + static_assert(dimensions >= 1 && dimensions <= 3, + "id can only be 1, 2, or 3 dimensional."); + template + using ParamTy = detail::enable_if_t<(N == val), T>; + + public: + id() = default; + +#ifdef __SYCL_DISABLE_ID_TO_INT_CONV__ + /* The following constructor is only available in the id struct + * specialization where: dimensions==1 */ + template id(ParamTy dim0) : base(dim0) {} + + template + id(ParamTy> &range_size) + : base(range_size.get(0)) {} + + template + id(ParamTy> &item) + : base(item.get_id(0)) {} +#endif // __SYCL_DISABLE_ID_TO_INT_CONV__ + + /* The following constructor is only available in the id struct + * specialization where: dimensions==2 */ + template + id(ParamTy dim0, size_t dim1) : base(dim0, dim1) {} + + template + id(ParamTy> &range_size) + : base(range_size.get(0), range_size.get(1)) {} + + template + id(ParamTy> &item) + : base(item.get_id(0), item.get_id(1)) {} + + /* The following constructor is only available in the id struct + * specialization where: dimensions==3 */ + template + id(ParamTy dim0, size_t dim1, size_t dim2) + : base(dim0, dim1, dim2) {} + + template + id(ParamTy> &range_size) + : base(range_size.get(0), range_size.get(1), range_size.get(2)) {} + + template + id(ParamTy> &item) + : base(item.get_id(0), item.get_id(1), item.get_id(2)) {} + + explicit operator range() const { + range result( + detail::InitializedVal::template get<0>()); + for (int i = 0; i < dimensions; ++i) { + result[i] = this->get(i); + } + return result; } - return result; - } - // OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >= - #define __SYCL_GEN_OPT(op) \ - id operator op(const id &rhs) const { \ - id result; \ - for (int i = 0; i < dimensions; ++i) { \ - result.common_array[i] = this->common_array[i] op rhs.common_array[i]; \ - } \ - return result; \ +// OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >= +#define __SYCL_GEN_OPT(op) \ + id operator op(const id &rhs) const { \ + id result; \ + for (int i = 0; i < dimensions; ++i) { \ + result.common_array[i] = this->common_array[i] op rhs.common_array[i]; \ } \ - id operator op(const size_t &rhs) const { \ - id result; \ - for (int i = 0; i < dimensions; ++i) { \ - result.common_array[i] = this->common_array[i] op rhs; \ - } \ - return result; \ + return result; \ + } \ + id operator op(const size_t &rhs) const { \ + id result; \ + for (int i = 0; i < dimensions; ++i) { \ + result.common_array[i] = this->common_array[i] op rhs; \ } \ - friend id operator op(const size_t &lhs, \ - const id &rhs) { \ - id result; \ - for (int i = 0; i < dimensions; ++i) { \ - result.common_array[i] = lhs op rhs.common_array[i]; \ - } \ - return result; \ + return result; \ + } \ + friend id operator op(const size_t &lhs, \ + const id &rhs) { \ + id result; \ + for (int i = 0; i < dimensions; ++i) { \ + result.common_array[i] = lhs op rhs.common_array[i]; \ } \ + return result; \ + } + __SYCL_GEN_OPT(+) + __SYCL_GEN_OPT(-) + __SYCL_GEN_OPT(*) + __SYCL_GEN_OPT(/) + __SYCL_GEN_OPT(%) + __SYCL_GEN_OPT(<<) + __SYCL_GEN_OPT(>>) + __SYCL_GEN_OPT(&) + __SYCL_GEN_OPT(|) + __SYCL_GEN_OPT(^) + __SYCL_GEN_OPT(&&) + __SYCL_GEN_OPT(||) + __SYCL_GEN_OPT(<) + __SYCL_GEN_OPT(>) + __SYCL_GEN_OPT(<=) + __SYCL_GEN_OPT(>=) - __SYCL_GEN_OPT(+) - __SYCL_GEN_OPT(-) - __SYCL_GEN_OPT(*) - __SYCL_GEN_OPT(/) - __SYCL_GEN_OPT(%) - __SYCL_GEN_OPT(<<) - __SYCL_GEN_OPT(>>) - __SYCL_GEN_OPT(&) - __SYCL_GEN_OPT(|) - __SYCL_GEN_OPT(^) - __SYCL_GEN_OPT(&&) - __SYCL_GEN_OPT(||) - __SYCL_GEN_OPT(<) - __SYCL_GEN_OPT(>) - __SYCL_GEN_OPT(<=) - __SYCL_GEN_OPT(>=) - - #undef __SYCL_GEN_OPT - - // OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^= - #define __SYCL_GEN_OPT(op) \ - id &operator op(const id &rhs) { \ - for (int i = 0; i < dimensions; ++i) { \ - this->common_array[i] op rhs.common_array[i]; \ - } \ - return *this; \ +#undef __SYCL_GEN_OPT + +// OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^= +#define __SYCL_GEN_OPT(op) \ + id &operator op(const id &rhs) { \ + for (int i = 0; i < dimensions; ++i) { \ + this->common_array[i] op rhs.common_array[i]; \ } \ - id &operator op(const size_t &rhs) { \ - for (int i = 0; i < dimensions; ++i) { \ - this->common_array[i] op rhs; \ - } \ - return *this; \ + return *this; \ + } \ + id &operator op(const size_t &rhs) { \ + for (int i = 0; i < dimensions; ++i) { \ + this->common_array[i] op rhs; \ } \ + return *this; \ + } + + __SYCL_GEN_OPT(+=) + __SYCL_GEN_OPT(-=) + __SYCL_GEN_OPT(*=) + __SYCL_GEN_OPT(/=) + __SYCL_GEN_OPT(%=) + __SYCL_GEN_OPT(<<=) + __SYCL_GEN_OPT(>>=) + __SYCL_GEN_OPT(&=) + __SYCL_GEN_OPT(|=) + __SYCL_GEN_OPT(^=) + +#undef __SYCL_GEN_OPT + }; - __SYCL_GEN_OPT(+=) - __SYCL_GEN_OPT(-=) - __SYCL_GEN_OPT(*=) - __SYCL_GEN_OPT(/=) - __SYCL_GEN_OPT(%=) - __SYCL_GEN_OPT(<<=) - __SYCL_GEN_OPT(>>=) - __SYCL_GEN_OPT(&=) - __SYCL_GEN_OPT(|=) - __SYCL_GEN_OPT(^=) - - #undef __SYCL_GEN_OPT -}; - -namespace detail { -template -size_t getOffsetForId(range Range, id Id, - id Offset) { - size_t offset = 0; - for (int i = 0; i < dimensions; ++i) - offset = offset * Range[i] + Offset[i] + Id[i]; - return offset; -} -} // namespace detail +#ifndef __SYCL_DISABLE_ID_TO_INT_CONV__ + template <> class id<1> : public detail::array<1> { + private: + using base = detail::array<1>; + + public: + id() = default; + + /* The following constructor is only available in the id struct + * specialization where: dimensions==1 */ + id<1>(size_t dim0) : base(dim0) {} + + id<1>(const range<1> &range_size) : base(range_size.get(0)) {} + + template + id<1>(const item<1, with_offset> &item) : base(item.get_id(0)) {} + + explicit operator range<1>() const { + range<1> result(detail::InitializedVal<1, range>::template get<0>()); + result[0] = this->get(0); + return result; + } + + operator size_t() const { return this->get(0); } + +// OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >= +#define __SYCL_GEN_OPT(op) \ + id<1> operator op(const id<1> &rhs) const { \ + id<1> result; \ + result.common_array[0] = this->common_array[0] op rhs.common_array[0]; \ + return result; \ + } \ + template id<1> operator op(const T &rhs) const { \ + id<1> result; \ + result.common_array[0] = this->common_array[0] op rhs; \ + return result; \ + } \ + template \ + friend id<1> operator op(const T &lhs, const id<1> &rhs) { \ + id<1> result; \ + result.common_array[0] = lhs op rhs.common_array[0]; \ + return result; \ + } \ + id<1> operator op(const range<1> &rhs) const { \ + id<1> result; \ + result.common_array[0] = this->common_array[0] op rhs[0]; \ + return result; \ + } \ + friend id<1> operator op(const range<1> &lhs, const id<1> &rhs) { \ + id<1> result; \ + result.common_array[0] = lhs[0] op rhs.common_array[0]; \ + return result; \ + } + + __SYCL_GEN_OPT(+) + __SYCL_GEN_OPT(-) + __SYCL_GEN_OPT(*) + __SYCL_GEN_OPT(/) + __SYCL_GEN_OPT(%) + __SYCL_GEN_OPT(<<) + __SYCL_GEN_OPT(>>) + __SYCL_GEN_OPT(&) + __SYCL_GEN_OPT(|) + __SYCL_GEN_OPT(^) + __SYCL_GEN_OPT(&&) + __SYCL_GEN_OPT(||) + __SYCL_GEN_OPT(<) + __SYCL_GEN_OPT(>) + __SYCL_GEN_OPT(<=) + __SYCL_GEN_OPT(>=) + +#undef __SYCL_GEN_OPT + +// OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^= +#define __SYCL_GEN_OPT(op) \ + id<1> &operator op(const id<1> &rhs) { \ + this->common_array[0] op rhs.common_array[0]; \ + return *this; \ + } \ + id<1> &operator op(const size_t &rhs) { \ + this->common_array[0] op rhs; \ + return *this; \ + } + + __SYCL_GEN_OPT(+=) + __SYCL_GEN_OPT(-=) + __SYCL_GEN_OPT(*=) + __SYCL_GEN_OPT(/=) + __SYCL_GEN_OPT(%=) + __SYCL_GEN_OPT(<<=) + __SYCL_GEN_OPT(>>=) + __SYCL_GEN_OPT(&=) + __SYCL_GEN_OPT(|=) + __SYCL_GEN_OPT(^=) + +#undef __SYCL_GEN_OPT + }; +#endif // __SYCL_DISABLE_ID_TO_INT_CONV__ + + namespace detail { + template + size_t getOffsetForId(range Range, id Id, + id Offset) { + size_t offset = 0; + for (int i = 0; i < dimensions; ++i) + offset = offset * Range[i] + Offset[i] + Id[i]; + return offset; + } + } // namespace detail // C++ feature test macros are supported by all supported compilers // with the exception of MSVC 1914. It doesn't support deduction guides. #ifdef __cpp_deduction_guides -id(size_t)->id<1>; -id(size_t, size_t)->id<2>; -id(size_t, size_t, size_t)->id<3>; + id(size_t)->id<1>; + id(size_t, size_t)->id<2>; + id(size_t, size_t, size_t)->id<3>; #endif -} // namespace sycl + } // namespace sycl } // namespace cl diff --git a/sycl/test/basic_tests/id.cpp b/sycl/test/basic_tests/id.cpp index e3ef20d3c48ca..c16e259c41e90 100644 --- a/sycl/test/basic_tests/id.cpp +++ b/sycl/test/basic_tests/id.cpp @@ -1,5 +1,7 @@ // RUN: %clangxx -fsycl %s -o %t.out // RUN: %t.out +// RUN: %clangxx -D__SYCL_DISABLE_ID_TO_INT_CONV__ -fsycl %s -o %t_dis.out +// RUN: %t_dis.out //==--------------- id.cpp - SYCL id test ----------------------------------==// // @@ -78,23 +80,201 @@ int main() { /* size_t &operator[](int dimension)const * Return a reference to the requested dimension of the id object. */ + cl::sycl::id<1> one_dim_id_brackets(64); + assert(one_dim_id_brackets[0] == 64); + cl::sycl::id<2> two_dim_id_brackets(128, 256); + assert(two_dim_id_brackets[0] == 128 && two_dim_id_brackets[1] == 256); + cl::sycl::id<3> three_dim_id_brackets(64, 1, 2); + assert(three_dim_id_brackets[0] == 64 && three_dim_id_brackets[1] == 1 && + three_dim_id_brackets[2] == 2); - /* size_t &operator[](int dimension)const - * Return a reference to the requested dimension of the id object. */ +/* size_t &operator[](int dimension)const + * Return a reference to the requested dimension of the id object. */ - /* id operatorOP(const id &rhs) const - * Where OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=. - * Constructs and returns a new instance of the SYCL id class template with - * the same dimensionality as this SYCL id, where each element of the new SYCL - * id instance is the result of an element-wise OP operator between each - * element of this SYCL id and each element of the rhs id. If the operator - * returns a bool the result is the cast to size_t */ +/* id operatorOP(const id &rhs) const + * Where OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=. + * Constructs and returns a new instance of the SYCL id class template with + * the same dimensionality as this SYCL id, where each element of the new SYCL + * id instance is the result of an element-wise OP operator between each + * element of this SYCL id and each element of the rhs id. If the operator + * returns a bool the result is the cast to size_t */ +#define oneLeftValue 10 +#define oneRightValue 2 +#define twoLeftValue 15 +#define twoRightValue 7 +#define threeLeftValue 3 +#define threeRightValue 9 + + cl::sycl::id<1> one_dim_op_left(oneLeftValue); + cl::sycl::id<1> one_dim_op_right(oneRightValue); + cl::sycl::range<1> one_dim_op_range(oneRightValue); + + cl::sycl::id<2> two_dim_op_left(oneLeftValue, twoLeftValue); + cl::sycl::id<2> two_dim_op_right(oneRightValue, twoRightValue); + cl::sycl::range<2> two_dim_op_range(oneRightValue, twoRightValue); + + cl::sycl::id<3> three_dim_op_left(oneLeftValue, twoLeftValue, threeLeftValue); + cl::sycl::id<3> three_dim_op_right(oneRightValue, twoRightValue, + threeRightValue); + cl::sycl::range<3> three_dim_op_range(oneRightValue, twoRightValue, + threeRightValue); +#define OPERATOR_TEST_BASIC(op) \ + assert((one_dim_op_left op one_dim_op_right)[0] == \ + (oneLeftValue op oneRightValue)); \ + assert((one_dim_op_right op one_dim_op_left)[0] == \ + (oneRightValue op oneLeftValue)); \ + assert((one_dim_op_left op oneRightValue)[0] == \ + (oneLeftValue op oneRightValue)); \ + assert((oneLeftValue op one_dim_op_right)[0] == \ + (oneLeftValue op oneRightValue)); \ + assert(((two_dim_op_left op two_dim_op_right)[0] == \ + (oneLeftValue op oneRightValue)) && \ + ((two_dim_op_right op two_dim_op_left)[1] == \ + (twoRightValue op twoLeftValue))); \ + assert(((two_dim_op_left op oneRightValue)[0] == \ + (oneLeftValue op oneRightValue)) && \ + ((twoLeftValue op two_dim_op_right)[1] == \ + (twoLeftValue op twoRightValue))); \ + assert(((three_dim_op_left op three_dim_op_right)[0] == \ + (oneLeftValue op oneRightValue)) && \ + ((three_dim_op_left op three_dim_op_right)[1] == \ + (twoLeftValue op twoRightValue)) && \ + ((three_dim_op_left op three_dim_op_right)[2] == \ + (threeLeftValue op threeRightValue))); \ + assert(((three_dim_op_left op oneRightValue)[0] == \ + (oneLeftValue op oneRightValue)) && \ + ((twoLeftValue op three_dim_op_right)[1] == \ + (twoLeftValue op twoRightValue)) && \ + ((three_dim_op_left op threeRightValue)[2] == \ + (threeLeftValue op threeRightValue))); + +#ifndef __SYCL_DISABLE_ID_TO_INT_CONV__ +#define OPERATOR_TEST(op) \ + OPERATOR_TEST_BASIC(op) \ + assert((one_dim_op_left op one_dim_op_range)[0] == \ + (oneLeftValue op oneRightValue)); \ + assert((one_dim_op_range op one_dim_op_left)[0] == \ + (oneRightValue op oneLeftValue)); +#else +#define OPERATOR_TEST(op) OPERATOR_TEST_BASIC(op) +#endif // __SYCL_DISABLE_ID_TO_INT_CONV__ + + OPERATOR_TEST(+) + OPERATOR_TEST(-) + OPERATOR_TEST(*) + OPERATOR_TEST(/) + OPERATOR_TEST(%) + OPERATOR_TEST(<<) + OPERATOR_TEST(>>) + OPERATOR_TEST(&) + OPERATOR_TEST(|) + OPERATOR_TEST(^) + OPERATOR_TEST(&&) + OPERATOR_TEST(||) + OPERATOR_TEST(<) + OPERATOR_TEST(>) + OPERATOR_TEST(<=) + OPERATOR_TEST(>=) + +#undef OPERATOR_TEST +#undef OPERATOR_TEST_BASIC /* id operatorOP(const id &rhs) const - * Where OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=. + * Where OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=. * Constructs and returns a new instance of the SYCL id class template with * the same dimensionality as this SYCL id, where each element of the new SYCL * id instance is the result of an element-wise OP operator between each * element of this SYCL id and each element of the rhs id. If the operator * returns a bool the result is the cast to size_t */ + +#define OPERATOR_TEST_BASIC(op) \ + one_dim_op_left[0] = oneLeftValue; \ + one_dim_op_right[0] = oneRightValue; \ + assert((one_dim_op_left op## = one_dim_op_right)[0] == \ + (oneLeftValue op oneRightValue)); \ + one_dim_op_left[0] = oneLeftValue; \ + assert((one_dim_op_left op## = oneRightValue)[0] == \ + (oneLeftValue op oneRightValue)); \ + two_dim_op_left[0] = oneLeftValue; \ + two_dim_op_left[1] = twoLeftValue; \ + two_dim_op_right[0] = oneRightValue; \ + two_dim_op_right[1] = twoRightValue; \ + assert(((two_dim_op_left op## = two_dim_op_right)[0] == \ + (oneLeftValue op oneRightValue)) && \ + (two_dim_op_left[1] == (twoLeftValue op twoRightValue))); \ + two_dim_op_left[0] = oneLeftValue; \ + two_dim_op_left[1] = twoLeftValue; \ + assert(((two_dim_op_left op## = oneRightValue)[0] == \ + (oneLeftValue op oneRightValue)) && \ + (two_dim_op_left[1] == (twoLeftValue op oneRightValue))); \ + three_dim_op_left[0] = oneLeftValue; \ + three_dim_op_left[1] = twoLeftValue; \ + three_dim_op_left[2] = threeLeftValue; \ + three_dim_op_right[0] = oneRightValue; \ + three_dim_op_right[1] = twoRightValue; \ + three_dim_op_right[2] = threeRightValue; \ + assert(((three_dim_op_left op## = three_dim_op_right)[0] == \ + (oneLeftValue op oneRightValue)) && \ + (three_dim_op_left[1] == (twoLeftValue op twoRightValue)) && \ + (three_dim_op_left[2] == (threeLeftValue op threeRightValue))); \ + three_dim_op_left[0] = oneLeftValue; \ + three_dim_op_left[1] = twoLeftValue; \ + three_dim_op_left[2] = threeLeftValue; \ + assert(((three_dim_op_left op## = oneRightValue)[0] == \ + (oneLeftValue op oneRightValue)) && \ + (three_dim_op_left[1] == (twoLeftValue op oneRightValue)) && \ + (three_dim_op_left[2] == (threeLeftValue op oneRightValue))); + +#ifndef __SYCL_DISABLE_ID_TO_INT_CONV__ +#define OPERATOR_TEST(op) \ + OPERATOR_TEST_BASIC(op) \ + one_dim_op_left[0] = oneLeftValue; \ + assert((one_dim_op_left op## = one_dim_op_range)[0] == \ + (oneLeftValue op oneRightValue)); \ + two_dim_op_left[0] = oneLeftValue; \ + two_dim_op_left[1] = twoLeftValue; \ + two_dim_op_range[0] = oneRightValue; \ + two_dim_op_range[1] = twoRightValue; +#else +#define OPERATOR_TEST(op) OPERATOR_TEST_BASIC(op) +#endif // __SYCL_DISABLE_ID_TO_INT_CONV__ + + OPERATOR_TEST(+) + OPERATOR_TEST(-) + OPERATOR_TEST(*) + OPERATOR_TEST(/) + OPERATOR_TEST(%) + OPERATOR_TEST(<<) + OPERATOR_TEST(>>) + OPERATOR_TEST(&) + OPERATOR_TEST(|) + OPERATOR_TEST(^) + +#undef OPERATOR_TEST +#undef OPERATOR_TEST_BASIC + +#undef oneLeftValue +#undef oneRightValue +#undef twoLeftValue +#undef twoRightValue +#undef threeLeftValue +#undef threeRightValue + +#ifndef __SYCL_DISABLE_ID_TO_INT_CONV__ +/* operator size_t() const + * Test implicit cast from id<1> to size_t and int value + * Should fails on cast from id<2> and id<3> */ +#define numValue 16 + cl::sycl::id<1> one_dim_id_cast_to_num(numValue); + size_t number1 = one_dim_id_cast_to_num; + int number2 = one_dim_id_cast_to_num; + size_t number3 = (size_t)one_dim_id_cast_to_num; + int number4 = (int)one_dim_id_cast_to_num; + size_t number5 = (int)one_dim_id_cast_to_num; + assert((number1 == numValue) && (number2 == numValue) && + (number3 == numValue) && (number4 == numValue) && + (number5 == numValue)); + +#undef numValue +#endif // __SYCL_DISABLE_ID_TO_INT_CONV__ }