diff --git a/sycl/include/sycl/h_item.hpp b/sycl/include/sycl/h_item.hpp index 63d79fc58de3a..8f9123e7ad734 100644 --- a/sycl/include/sycl/h_item.hpp +++ b/sycl/include/sycl/h_item.hpp @@ -25,8 +25,10 @@ class Builder; /// call or to the corresponding parallel_for_work_group call. /// /// \ingroup sycl_api -template class h_item { +template class h_item { public: + static constexpr int dimensions = Dimensions; + h_item() = delete; h_item(const h_item &hi) = default; @@ -34,70 +36,70 @@ template class h_item { h_item &operator=(const h_item &hi) = default; /* -- public interface members -- */ - item get_global() const { return globalItem; } + item get_global() const { return globalItem; } - item get_local() const { return get_logical_local(); } + item get_local() const { return get_logical_local(); } - item get_logical_local() const { return logicalLocalItem; } + item get_logical_local() const { return logicalLocalItem; } - item get_physical_local() const { return localItem; } + item get_physical_local() const { return localItem; } - range get_global_range() const { + range get_global_range() const { return get_global().get_range(); } - size_t get_global_range(int dimension) const { - return get_global().get_range(dimension); + size_t get_global_range(int Dimension) const { + return get_global().get_range(Dimension); } - id get_global_id() const { return get_global().get_id(); } + id get_global_id() const { return get_global().get_id(); } - size_t get_global_id(int dimension) const { - return get_global().get_id(dimension); + size_t get_global_id(int Dimension) const { + return get_global().get_id(Dimension); } - range get_local_range() const { return get_local().get_range(); } + range get_local_range() const { return get_local().get_range(); } - size_t get_local_range(int dimension) const { - return get_local().get_range(dimension); + size_t get_local_range(int Dimension) const { + return get_local().get_range(Dimension); } - id get_local_id() const { return get_local().get_id(); } + id get_local_id() const { return get_local().get_id(); } - size_t get_local_id(int dimension) const { - return get_local().get_id(dimension); + size_t get_local_id(int Dimension) const { + return get_local().get_id(Dimension); } - range get_logical_local_range() const { + range get_logical_local_range() const { return get_logical_local().get_range(); } - size_t get_logical_local_range(int dimension) const { - return get_logical_local().get_range(dimension); + size_t get_logical_local_range(int Dimension) const { + return get_logical_local().get_range(Dimension); } - id get_logical_local_id() const { + id get_logical_local_id() const { return get_logical_local().get_id(); } - size_t get_logical_local_id(int dimension) const { - return get_logical_local().get_id(dimension); + size_t get_logical_local_id(int Dimension) const { + return get_logical_local().get_id(Dimension); } - range get_physical_local_range() const { + range get_physical_local_range() const { return get_physical_local().get_range(); } - size_t get_physical_local_range(int dimension) const { - return get_physical_local().get_range(dimension); + size_t get_physical_local_range(int Dimension) const { + return get_physical_local().get_range(Dimension); } - id get_physical_local_id() const { + id get_physical_local_id() const { return get_physical_local().get_id(); } - size_t get_physical_local_id(int dimension) const { - return get_physical_local().get_id(dimension); + size_t get_physical_local_id(int Dimension) const { + return get_physical_local().get_id(Dimension); } bool operator==(const h_item &rhs) const { @@ -109,29 +111,29 @@ template class h_item { protected: friend class detail::Builder; - friend class group; - h_item(const item &GL, const item &L, - const range &flexLocalRange) + friend class group; + h_item(const item &GL, const item &L, + const range &flexLocalRange) : globalItem(GL), localItem(L), - logicalLocalItem(detail::Builder::createItem( + logicalLocalItem(detail::Builder::createItem( flexLocalRange, L.get_id())) {} - h_item(const item &GL, const item &L) + h_item(const item &GL, const item &L) : globalItem(GL), localItem(L), - logicalLocalItem(detail::Builder::createItem( + logicalLocalItem(detail::Builder::createItem( localItem.get_range(), localItem.get_id())) {} - void setLogicalLocalID(const id &ID) { + void setLogicalLocalID(const id &ID) { detail::Builder::updateItemIndex(logicalLocalItem, ID); } private: /// Describles physical workgroup range and current \c h_item position in it. - item localItem; + item localItem; /// Describles global range and current \c h_item position in it. - item globalItem; + item globalItem; /// Describles logical flexible range and current \c h_item position in it. - item logicalLocalItem; + item logicalLocalItem; }; } // __SYCL_INLINE_VER_NAMESPACE(_V1) diff --git a/sycl/include/sycl/id.hpp b/sycl/include/sycl/id.hpp index a71a3f785b947..8aa45a528089f 100644 --- a/sycl/include/sycl/id.hpp +++ b/sycl/include/sycl/id.hpp @@ -23,17 +23,20 @@ class RoundedRangeKernel; template class RoundedRangeKernelWithKH; } // namespace detail -template class range; -template class item; +template class range; +template class item; /// A unique identifier of an item in an index space. /// /// \ingroup sycl_api -template class id : public detail::array { +template class id : public detail::array { +public: + static constexpr int dimensions = Dimensions; + private: - using base = detail::array; - static_assert(dimensions >= 1 && dimensions <= 3, - "id can only be 1, 2, or 3 dimensional."); + using base = detail::array; + static_assert(Dimensions >= 1 && Dimensions <= 3, + "id can only be 1, 2, or 3 Dimensional."); template using ParamTy = std::enable_if_t<(N == val), T>; @@ -54,49 +57,49 @@ template class id : public detail::array { id() = default; /* The following constructor is only available in the id struct - * specialization where: dimensions==1 */ - template id(ParamTy dim0) : base(dim0) {} + * specialization where: Dimensions==1 */ + template id(ParamTy dim0) : base(dim0) {} - template - id(ParamTy> &range_size) + template + id(ParamTy> &range_size) : base(range_size.get(0)) {} - template - id(ParamTy> &item) + template + id(ParamTy> &item) : base(item.get_id(0)) {} /* The following constructor is only available in the id struct - * specialization where: dimensions==2 */ - template + * specialization where: Dimensions==2 */ + template id(ParamTy dim0, size_t dim1) : base(dim0, dim1) {} - template - id(ParamTy> &range_size) + template + id(ParamTy> &range_size) : base(range_size.get(0), range_size.get(1)) {} - template - id(ParamTy> &item) + 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 + * specialization where: Dimensions==3 */ + template id(ParamTy dim0, size_t dim1, size_t dim2) : base(dim0, dim1, dim2) {} - template - id(ParamTy> &range_size) + template + id(ParamTy> &range_size) : base(range_size.get(0), range_size.get(1), range_size.get(2)) {} - template - id(ParamTy> &item) + template + id(ParamTy> &item) : base(item.get_id(0), item.get_id(1), item.get_id(2)) {} __SYCL_DEPRECATED("range() conversion is deprecated") - explicit operator range() const { - range result( - detail::InitializedVal::template get<0>()); - for (int i = 0; i < dimensions; ++i) { + 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; @@ -108,7 +111,7 @@ template class id : public detail::array { * conversion: * int a = id<1>(value); */ - __SYCL_ALWAYS_INLINE operator EnableIfT<(dimensions == 1), size_t>() const { + __SYCL_ALWAYS_INLINE operator EnableIfT<(Dimensions == 1), size_t>() const { size_t Result = this->common_array[0]; __SYCL_ASSUME_INT(Result); return Result; @@ -117,14 +120,14 @@ template class id : public detail::array { // OP is: ==, != #ifndef __SYCL_DISABLE_ID_TO_INT_CONV__ - using detail::array::operator==; + using detail::array::operator==; // Needed for clang in C++20 mode as the above operator== would be ambigious // between regular/reversed call for "Id == Id" case. - bool operator==(const id &rhs) const { - return this->detail::array::operator==(rhs); + bool operator==(const id &rhs) const { + return this->detail::array::operator==(rhs); } #if __cpp_impl_three_way_comparison < 201907 - using detail::array::operator!=; + using detail::array::operator!=; #endif /* Enable operators with integral types. @@ -141,7 +144,7 @@ template class id : public detail::array { } \ template \ friend EnableIfIntegral operator op(const T &lhs, \ - const id &rhs) { \ + const id &rhs) { \ if (lhs != rhs.common_array[0]) \ return false op true; \ return true op true; \ @@ -156,10 +159,10 @@ template class id : public detail::array { // OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >= #define __SYCL_GEN_OPT_BASE(op) \ - friend id operator op(const id &lhs, \ - const id &rhs) { \ - id result; \ - for (int i = 0; i < dimensions; ++i) { \ + friend id operator op(const id &lhs, \ + const id &rhs) { \ + id result; \ + for (int i = 0; i < Dimensions; ++i) { \ result.common_array[i] = lhs.common_array[i] op rhs.common_array[i]; \ } \ return result; \ @@ -170,19 +173,19 @@ template class id : public detail::array { #define __SYCL_GEN_OPT(op) \ __SYCL_GEN_OPT_BASE(op) \ template \ - friend EnableIfIntegral> operator op( \ - const id &lhs, const T &rhs) { \ - id result; \ - for (int i = 0; i < dimensions; ++i) { \ + friend EnableIfIntegral> operator op( \ + const id &lhs, const T &rhs) { \ + id result; \ + for (int i = 0; i < Dimensions; ++i) { \ result.common_array[i] = lhs.common_array[i] op rhs; \ } \ return result; \ } \ template \ - friend EnableIfIntegral> operator op( \ - const T &lhs, const id &rhs) { \ - id result; \ - for (int i = 0; i < dimensions; ++i) { \ + friend EnableIfIntegral> operator op( \ + const 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; \ @@ -190,18 +193,18 @@ template class id : public detail::array { #else #define __SYCL_GEN_OPT(op) \ __SYCL_GEN_OPT_BASE(op) \ - friend id operator op(const id &lhs, \ + friend id operator op(const id &lhs, \ const size_t &rhs) { \ - id result; \ - for (int i = 0; i < dimensions; ++i) { \ + id result; \ + for (int i = 0; i < Dimensions; ++i) { \ result.common_array[i] = lhs.common_array[i] op rhs; \ } \ return result; \ } \ - friend id operator op(const size_t &lhs, \ - const id &rhs) { \ - id result; \ - for (int i = 0; i < dimensions; ++i) { \ + 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; \ @@ -230,15 +233,15 @@ template class id : public detail::array { // OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^= #define __SYCL_GEN_OPT(op) \ - friend id &operator op(id &lhs, \ - const id &rhs) { \ - for (int i = 0; i < dimensions; ++i) { \ + friend id &operator op(id &lhs, \ + const id &rhs) { \ + for (int i = 0; i < Dimensions; ++i) { \ lhs.common_array[i] op rhs.common_array[i]; \ } \ return lhs; \ } \ - friend id &operator op(id &lhs, const size_t &rhs) { \ - for (int i = 0; i < dimensions; ++i) { \ + friend id &operator op(id &lhs, const size_t &rhs) { \ + for (int i = 0; i < Dimensions; ++i) { \ lhs.common_array[i] op rhs; \ } \ return lhs; \ @@ -259,9 +262,9 @@ template class id : public detail::array { // OP is unary +, - #define __SYCL_GEN_OPT(op) \ - friend id operator op(const id &rhs) { \ - id result; \ - for (int i = 0; i < dimensions; ++i) { \ + friend id operator op(const id &rhs) { \ + id result; \ + for (int i = 0; i < Dimensions; ++i) { \ result.common_array[i] = (op rhs.common_array[i]); \ } \ return result; \ @@ -274,8 +277,8 @@ template class id : public detail::array { // OP is prefix ++, -- #define __SYCL_GEN_OPT(op) \ - friend id &operator op(id &rhs) { \ - for (int i = 0; i < dimensions; ++i) { \ + friend id &operator op(id &rhs) { \ + for (int i = 0; i < Dimensions; ++i) { \ op rhs.common_array[i]; \ } \ return rhs; \ @@ -288,9 +291,9 @@ template class id : public detail::array { // OP is postfix ++, -- #define __SYCL_GEN_OPT(op) \ - friend id operator op(id &lhs, int) { \ - id old_lhs; \ - for (int i = 0; i < dimensions; ++i) { \ + friend id operator op(id &lhs, int) { \ + id old_lhs; \ + for (int i = 0; i < Dimensions; ++i) { \ old_lhs.common_array[i] = lhs.common_array[i]; \ op lhs.common_array[i]; \ } \ @@ -307,15 +310,15 @@ template class id : public detail::array { template friend class detail::RoundedRangeKernel; template friend class detail::RoundedRangeKernelWithKH; - void set_allowed_range(range rnwi) { (void)rnwi[0]; } + void set_allowed_range(range rnwi) { (void)rnwi[0]; } }; namespace detail { -template -size_t getOffsetForId(range Range, id Id, - id Offset) { +template +size_t getOffsetForId(range Range, id Id, + id Offset) { size_t offset = 0; - for (int i = 0; i < dimensions; ++i) + for (int i = 0; i < Dimensions; ++i) offset = offset * Range[i] + Offset[i] + Id[i]; return offset; } diff --git a/sycl/include/sycl/item.hpp b/sycl/include/sycl/item.hpp index fc0f161bee62a..54c0207b35128 100644 --- a/sycl/include/sycl/item.hpp +++ b/sycl/include/sycl/item.hpp @@ -36,7 +36,11 @@ item getDelinearizedItem(range Range, id Id); /// in a range. /// /// \ingroup sycl_api -template class item { +template class item { +public: + static constexpr int dimensions = Dimensions; + +private: #ifndef __SYCL_DISABLE_ITEM_TO_INT_CONV__ /* Helper class for conversion operator. Void type is not suitable. User * cannot even try to get address of the operator __private_class(). User @@ -50,48 +54,48 @@ template class item { public: item() = delete; - id get_id() const { return MImpl.MIndex; } + id get_id() const { return MImpl.MIndex; } - size_t __SYCL_ALWAYS_INLINE get_id(int dimension) const { - size_t Id = MImpl.MIndex[dimension]; + size_t __SYCL_ALWAYS_INLINE get_id(int Dimension) const { + size_t Id = MImpl.MIndex[Dimension]; __SYCL_ASSUME_INT(Id); return Id; } - size_t __SYCL_ALWAYS_INLINE operator[](int dimension) const { - size_t Id = MImpl.MIndex[dimension]; + size_t __SYCL_ALWAYS_INLINE operator[](int Dimension) const { + size_t Id = MImpl.MIndex[Dimension]; __SYCL_ASSUME_INT(Id); return Id; } - range get_range() const { return MImpl.MExtent; } + range get_range() const { return MImpl.MExtent; } - size_t __SYCL_ALWAYS_INLINE get_range(int dimension) const { - size_t Id = MImpl.MExtent[dimension]; + size_t __SYCL_ALWAYS_INLINE get_range(int Dimension) const { + size_t Id = MImpl.MExtent[Dimension]; __SYCL_ASSUME_INT(Id); return Id; } #ifndef __SYCL_DISABLE_ITEM_TO_INT_CONV__ - operator EnableIfT() const { return get_id(0); } + operator EnableIfT() const { return get_id(0); } #endif // __SYCL_DISABLE_ITEM_TO_INT_CONV__ template __SYCL2020_DEPRECATED("offsets are deprecated in SYCL2020") - std::enable_if_t> get_offset() const { + std::enable_if_t> get_offset() const { return MImpl.MOffset; } template __SYCL2020_DEPRECATED("offsets are deprecated in SYCL2020") std::enable_if_t __SYCL_ALWAYS_INLINE - get_offset(int dimension) const { - size_t Id = MImpl.MOffset[dimension]; + get_offset(int Dimension) const { + size_t Id = MImpl.MOffset[Dimension]; __SYCL_ASSUME_INT(Id); return Id; } template - operator std::enable_if_t>() const { - return detail::Builder::createItem( + operator std::enable_if_t>() const { + return detail::Builder::createItem( MImpl.MExtent, MImpl.MIndex, /*Offset*/ {}); } @@ -103,7 +107,7 @@ template class item { item(const item &rhs) = default; - item(item &&rhs) = default; + item(item &&rhs) = default; item &operator=(const item &rhs) = default; @@ -115,13 +119,13 @@ template class item { protected: template - item(std::enable_if_t> &extent, - const id &index, const id &offset) + item(std::enable_if_t> &extent, + const id &index, const id &offset) : MImpl{extent, index, offset} {} template - item(std::enable_if_t> &extent, - const id &index) + item(std::enable_if_t> &extent, + const id &index) : MImpl{extent, index} {} friend class detail::Builder; @@ -131,13 +135,13 @@ template class item { template friend class detail::RoundedRangeKernel; template friend class detail::RoundedRangeKernelWithKH; - void set_allowed_range(const range rnwi) { MImpl.MExtent = rnwi; } + void set_allowed_range(const range rnwi) { MImpl.MExtent = rnwi; } template friend item detail::reduction::getDelinearizedItem(range Range, id Id); - detail::ItemBase MImpl; + detail::ItemBase MImpl; }; template diff --git a/sycl/include/sycl/nd_item.hpp b/sycl/include/sycl/nd_item.hpp index f55379309acd8..9b53a0addd65a 100644 --- a/sycl/include/sycl/nd_item.hpp +++ b/sycl/include/sycl/nd_item.hpp @@ -30,21 +30,23 @@ class Builder; } namespace ext::oneapi::experimental { -template class root_group; +template class root_group; } /// Identifies an instance of the function object executing at each point in an /// nd_range. /// /// \ingroup sycl_api -template class nd_item { +template class nd_item { public: + static constexpr int dimensions = Dimensions; + nd_item() = delete; - id get_global_id() const { return globalItem.get_id(); } + id get_global_id() const { return globalItem.get_id(); } - size_t __SYCL_ALWAYS_INLINE get_global_id(int dimension) const { - size_t Id = globalItem.get_id(dimension); + size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const { + size_t Id = globalItem.get_id(Dimension); __SYCL_ASSUME_INT(Id); return Id; } @@ -55,10 +57,10 @@ template class nd_item { return Id; } - id get_local_id() const { return localItem.get_id(); } + id get_local_id() const { return localItem.get_id(); } - size_t __SYCL_ALWAYS_INLINE get_local_id(int dimension) const { - size_t Id = localItem.get_id(dimension); + size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const { + size_t Id = localItem.get_id(Dimension); __SYCL_ASSUME_INT(Id); return Id; } @@ -69,12 +71,12 @@ template class nd_item { return Id; } - group get_group() const { return Group; } + group get_group() const { return Group; } sub_group get_sub_group() const { return sub_group(); } - size_t __SYCL_ALWAYS_INLINE get_group(int dimension) const { - size_t Size = Group[dimension]; + size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const { + size_t Size = Group[Dimension]; __SYCL_ASSUME_INT(Size); return Size; } @@ -85,31 +87,31 @@ template class nd_item { return Id; } - range get_group_range() const { return Group.get_group_range(); } + range get_group_range() const { return Group.get_group_range(); } - size_t __SYCL_ALWAYS_INLINE get_group_range(int dimension) const { - size_t Range = Group.get_group_range(dimension); + size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const { + size_t Range = Group.get_group_range(Dimension); __SYCL_ASSUME_INT(Range); return Range; } - range get_global_range() const { return globalItem.get_range(); } + range get_global_range() const { return globalItem.get_range(); } - size_t get_global_range(int dimension) const { - return globalItem.get_range(dimension); + size_t get_global_range(int Dimension) const { + return globalItem.get_range(Dimension); } - range get_local_range() const { return localItem.get_range(); } + range get_local_range() const { return localItem.get_range(); } - size_t get_local_range(int dimension) const { - return localItem.get_range(dimension); + size_t get_local_range(int Dimension) const { + return localItem.get_range(Dimension); } __SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020") - id get_offset() const { return globalItem.get_offset(); } + id get_offset() const { return globalItem.get_offset(); } - nd_range get_nd_range() const { - return nd_range(get_global_range(), get_local_range(), + nd_range get_nd_range() const { + return nd_range(get_global_range(), get_local_range(), get_offset()); } @@ -202,9 +204,9 @@ template class nd_item { Group.wait_for(events...); } - sycl::ext::oneapi::experimental::root_group + sycl::ext::oneapi::experimental::root_group ext_oneapi_get_root_group() const { - return sycl::ext::oneapi::experimental::root_group{*this}; + return sycl::ext::oneapi::experimental::root_group{*this}; } nd_item(const nd_item &rhs) = default; @@ -224,14 +226,14 @@ template class nd_item { protected: friend class detail::Builder; - nd_item(const item &GL, const item &L, - const group &GR) + nd_item(const item &GL, const item &L, + const group &GR) : globalItem(GL), localItem(L), Group(GR) {} private: - item globalItem; - item localItem; - group Group; + item globalItem; + item localItem; + group Group; }; template diff --git a/sycl/include/sycl/nd_range.hpp b/sycl/include/sycl/nd_range.hpp index 6d2d05a563186..4bb7ff97c959c 100644 --- a/sycl/include/sycl/nd_range.hpp +++ b/sycl/include/sycl/nd_range.hpp @@ -20,46 +20,50 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) { /// dispatch. /// /// \ingroup sycl_api -template class nd_range { - range globalSize; - range localSize; - id offset; - static_assert(dimensions >= 1 && dimensions <= 3, - "nd_range can only be 1, 2, or 3 dimensional."); +template class nd_range { +public: + static constexpr int dimensions = Dimensions; + +private: + range globalSize; + range localSize; + id offset; + static_assert(Dimensions >= 1 && Dimensions <= 3, + "nd_range can only be 1, 2, or 3 Dimensional."); public: __SYCL2020_DEPRECATED("offsets are deprecated in SYCL2020") - nd_range(range globalSize, range localSize, - id offset) + nd_range(range globalSize, range localSize, + id offset) : globalSize(globalSize), localSize(localSize), offset(offset) {} - nd_range(range globalSize, range localSize) - : globalSize(globalSize), localSize(localSize), offset(id()) { + nd_range(range globalSize, range localSize) + : globalSize(globalSize), localSize(localSize), offset(id()) { } - range get_global_range() const { return globalSize; } + range get_global_range() const { return globalSize; } - range get_local_range() const { return localSize; } + range get_local_range() const { return localSize; } - range get_group_range() const { return globalSize / localSize; } + range get_group_range() const { return globalSize / localSize; } __SYCL2020_DEPRECATED("offsets are deprecated in SYCL2020") - id get_offset() const { return offset; } + id get_offset() const { return offset; } // Common special member functions for by-value semantics - nd_range(const nd_range &rhs) = default; - nd_range(nd_range &&rhs) = default; - nd_range &operator=(const nd_range &rhs) = default; - nd_range &operator=(nd_range &&rhs) = default; + nd_range(const nd_range &rhs) = default; + nd_range(nd_range &&rhs) = default; + nd_range &operator=(const nd_range &rhs) = default; + nd_range &operator=(nd_range &&rhs) = default; nd_range() = default; // Common member functions for by-value semantics - bool operator==(const nd_range &rhs) const { + bool operator==(const nd_range &rhs) const { return (rhs.globalSize == this->globalSize) && (rhs.localSize == this->localSize) && (rhs.offset == this->offset); } - bool operator!=(const nd_range &rhs) const { + bool operator!=(const nd_range &rhs) const { return !(*this == rhs); } }; diff --git a/sycl/include/sycl/range.hpp b/sycl/include/sycl/range.hpp index f6d6868dd66af..7c3d367114046 100644 --- a/sycl/include/sycl/range.hpp +++ b/sycl/include/sycl/range.hpp @@ -15,58 +15,62 @@ namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { -template class id; +template class id; /// Defines the iteration domain of either a single work-group in a parallel -/// dispatch, or the overall dimensions of the dispatch. +/// dispatch, or the overall Dimensions of the dispatch. /// /// \ingroup sycl_api -template class range : public detail::array { - static_assert(dimensions >= 1 && dimensions <= 3, - "range can only be 1, 2, or 3 dimensional."); - using base = detail::array; +template class range : public detail::array { +public: + static constexpr int dimensions = Dimensions; + +private: + static_assert(Dimensions >= 1 && Dimensions <= 3, + "range can only be 1, 2, or 3 Dimensional."); + using base = detail::array; template using IntegralType = std::enable_if_t, T>; public: /* The following constructor is only available in the range class - specialization where: dimensions==1 */ - template + specialization where: Dimensions==1 */ + template range(typename std::enable_if_t<(N == 1), size_t> dim0) : base(dim0) {} /* The following constructor is only available in the range class - specialization where: dimensions==2 */ - template + specialization where: Dimensions==2 */ + template range(typename std::enable_if_t<(N == 2), size_t> dim0, size_t dim1) : base(dim0, dim1) {} /* The following constructor is only available in the range class - specialization where: dimensions==3 */ - template + specialization where: Dimensions==3 */ + template range(typename std::enable_if_t<(N == 3), size_t> dim0, size_t dim1, size_t dim2) : base(dim0, dim1, dim2) {} size_t size() const { size_t size = 1; - for (int i = 0; i < dimensions; ++i) { + for (int i = 0; i < Dimensions; ++i) { size *= this->get(i); } return size; } - range(const range &rhs) = default; - range(range &&rhs) = default; - range &operator=(const range &rhs) = default; - range &operator=(range &&rhs) = default; + range(const range &rhs) = default; + range(range &&rhs) = default; + range &operator=(const range &rhs) = default; + range &operator=(range &&rhs) = default; range() = delete; // OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >= #define __SYCL_GEN_OPT_BASE(op) \ - friend range operator op(const range &lhs, \ - const range &rhs) { \ - range result(lhs); \ - for (int i = 0; i < dimensions; ++i) { \ + friend range operator op(const range &lhs, \ + const range &rhs) { \ + range result(lhs); \ + for (int i = 0; i < Dimensions; ++i) { \ result.common_array[i] = lhs.common_array[i] op rhs.common_array[i]; \ } \ return result; \ @@ -77,19 +81,19 @@ template class range : public detail::array { #define __SYCL_GEN_OPT(op) \ __SYCL_GEN_OPT_BASE(op) \ template \ - friend IntegralType> operator op( \ - const range &lhs, const T &rhs) { \ - range result(lhs); \ - for (int i = 0; i < dimensions; ++i) { \ + friend IntegralType> operator op( \ + const range &lhs, const T &rhs) { \ + range result(lhs); \ + for (int i = 0; i < Dimensions; ++i) { \ result.common_array[i] = lhs.common_array[i] op rhs; \ } \ return result; \ } \ template \ - friend IntegralType> operator op( \ - const T &lhs, const range &rhs) { \ - range result(rhs); \ - for (int i = 0; i < dimensions; ++i) { \ + friend IntegralType> operator op( \ + const T &lhs, const range &rhs) { \ + range result(rhs); \ + for (int i = 0; i < Dimensions; ++i) { \ result.common_array[i] = lhs op rhs.common_array[i]; \ } \ return result; \ @@ -97,18 +101,18 @@ template class range : public detail::array { #else #define __SYCL_GEN_OPT(op) \ __SYCL_GEN_OPT_BASE(op) \ - friend range operator op(const range &lhs, \ + friend range operator op(const range &lhs, \ const size_t &rhs) { \ - range result(lhs); \ - for (int i = 0; i < dimensions; ++i) { \ + range result(lhs); \ + for (int i = 0; i < Dimensions; ++i) { \ result.common_array[i] = lhs.common_array[i] op rhs; \ } \ return result; \ } \ - friend range operator op(const size_t &lhs, \ - const range &rhs) { \ - range result(rhs); \ - for (int i = 0; i < dimensions; ++i) { \ + friend range operator op(const size_t &lhs, \ + const range &rhs) { \ + range result(rhs); \ + for (int i = 0; i < Dimensions; ++i) { \ result.common_array[i] = lhs op rhs.common_array[i]; \ } \ return result; \ @@ -137,16 +141,16 @@ template class range : public detail::array { // OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^= #define __SYCL_GEN_OPT(op) \ - friend range &operator op(range &lhs, \ - const range &rhs) { \ - for (int i = 0; i < dimensions; ++i) { \ + friend range &operator op(range &lhs, \ + const range &rhs) { \ + for (int i = 0; i < Dimensions; ++i) { \ lhs.common_array[i] op rhs[i]; \ } \ return lhs; \ } \ - friend range &operator op(range &lhs, \ + friend range &operator op(range &lhs, \ const size_t &rhs) { \ - for (int i = 0; i < dimensions; ++i) { \ + for (int i = 0; i < Dimensions; ++i) { \ lhs.common_array[i] op rhs; \ } \ return lhs; \ @@ -167,9 +171,9 @@ template class range : public detail::array { // OP is unary +, - #define __SYCL_GEN_OPT(op) \ - friend range operator op(const range &rhs) { \ - range result(rhs); \ - for (int i = 0; i < dimensions; ++i) { \ + friend range operator op(const range &rhs) { \ + range result(rhs); \ + for (int i = 0; i < Dimensions; ++i) { \ result.common_array[i] = (op rhs.common_array[i]); \ } \ return result; \ @@ -182,8 +186,8 @@ template class range : public detail::array { // OP is prefix ++, -- #define __SYCL_GEN_OPT(op) \ - friend range &operator op(range &rhs) { \ - for (int i = 0; i < dimensions; ++i) { \ + friend range &operator op(range &rhs) { \ + for (int i = 0; i < Dimensions; ++i) { \ op rhs.common_array[i]; \ } \ return rhs; \ @@ -196,9 +200,9 @@ template class range : public detail::array { // OP is postfix ++, -- #define __SYCL_GEN_OPT(op) \ - friend range operator op(range &lhs, int) { \ - range old_lhs(lhs); \ - for (int i = 0; i < dimensions; ++i) { \ + friend range operator op(range &lhs, int) { \ + range old_lhs(lhs); \ + for (int i = 0; i < Dimensions; ++i) { \ op lhs.common_array[i]; \ } \ return old_lhs; \ diff --git a/sycl/test/basic_tests/dimension.cpp b/sycl/test/basic_tests/dimension.cpp new file mode 100644 index 0000000000000..278789d1c6064 --- /dev/null +++ b/sycl/test/basic_tests/dimension.cpp @@ -0,0 +1,19 @@ +// RUN: %clangxx -fsycl -fsyntax-only %s -o %t.out +#include + +template void check_dim_member() { + using namespace sycl; + static_assert(range::dimensions == N); + static_assert(nd_range::dimensions == N); + static_assert(id::dimensions == N); + static_assert(item::dimensions == N); + static_assert(nd_item::dimensions == N); + static_assert(h_item::dimensions == N); + static_assert(group::dimensions == N); +} + +int main() { + check_dim_member<1>(); + check_dim_member<2>(); + check_dim_member<3>(); +} diff --git a/sycl/test/gdb/accessors.cpp b/sycl/test/gdb/accessors.cpp index 14d59727bd93a..b998b8de39355 100644 --- a/sycl/test/gdb/accessors.cpp +++ b/sycl/test/gdb/accessors.cpp @@ -23,7 +23,7 @@ void foo(sycl::buffer &BufA) { // CHECK: CXXMethodDecl {{.*}}getSize // CHECK: CXXMethodDecl {{.*}}getPtr -// CHECK-DEBUG-INFO: !DICompositeType(tag: DW_TAG_class_type, name: "accessor >", {{.*}}, templateParams: ![[TEMPL_METADATA:[0-9]+]] +// CHECK-DEBUG-INFO-DAG: !DICompositeType(tag: DW_TAG_class_type, name: "accessor >", {{.*}}, templateParams: ![[TEMPL_METADATA:[0-9]+]] // CHECK-DEBUG-INFO-DAG: ![[TEMPL_METADATA]] = !{![[DATA_T:[0-9]+]], ![[Dims:[0-9]+]], ![[AccMode:[0-9]+]], ![[AccTarget:[0-9]+]], ![[IsPlh:[0-9]+]], ![[PropListT:[0-9]+]]} // CHECK-DEBUG-INFO-DAG: ![[DATA_T]] = !DITemplateTypeParameter(name: "DataT" // CHECK-DEBUG-INFO-DAG: ![[Dims]] = !DITemplateValueParameter(name: "Dimensions"