From f20318c72276a5892f29bba54eb752c9a0e98f5c Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 13 Sep 2022 05:56:38 -0400 Subject: [PATCH 01/21] [SYCL] Implement accessor iterator Added base type for implementing accessor iterators. Implemented `accessor::begin()` and `accessor::end()` methods. --- sycl/include/sycl/accessor.hpp | 17 + .../include/sycl/detail/accessor_iterator.hpp | 437 ++++++++++++++++++ sycl/test/basic_tests/accessor/iterator.cpp | 35 ++ sycl/unittests/CMakeLists.txt | 1 + sycl/unittests/accessor/AccessorIterator.cpp | 232 ++++++++++ sycl/unittests/accessor/CMakeLists.txt | 3 + 6 files changed, 725 insertions(+) create mode 100644 sycl/include/sycl/detail/accessor_iterator.hpp create mode 100644 sycl/test/basic_tests/accessor/iterator.cpp create mode 100644 sycl/unittests/accessor/AccessorIterator.cpp create mode 100644 sycl/unittests/accessor/CMakeLists.txt diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index e6cf64c25a0d8..af2a580d6ea6c 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -12,6 +12,7 @@ #include #include #include +#include #include #include #include @@ -32,6 +33,7 @@ #include #include +#include #include @@ -972,6 +974,13 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : using reference = DataT &; using const_reference = const DataT &; + using iterator = + typename detail::__accessor_iterator; + using difference_type = + typename std::iterator_traits::difference_type; + // The list of accessor constructors with their arguments // -------+---------+-------+----+-----+-------------- // Dimensions = 0 @@ -1835,6 +1844,14 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; } bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); } + iterator begin() noexcept { + return iterator::__get_begin(this, get_offset(), get_range()); + } + + iterator end() noexcept { + return iterator::__get_end(this, get_offset(), get_range()); + } + private: #ifdef __SYCL_DEVICE_ONLY__ size_t getTotalOffset() const { diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp new file mode 100644 index 0000000000000..1e868f0783e6f --- /dev/null +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -0,0 +1,437 @@ +//==------------ accessor_iterator.hpp - SYCL standard header file ---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include + +#include +#include + +/// \file accessor_iterator.hpp +/// The file contains implementation of accessor iterator class. +/// +/// The reason why we can't use a plain pointer as an interator and have to +/// implement a custom class here is explained in section 4.7.6.8. Ranged +/// accessors of SYCL 2020 specification. A couple of quotes from there: +/// +/// > Accessors of type accessor and host_accessor can be constructed from a +/// > sub-range of a buffer by providing a range and offset to the constructor. +/// > +/// > If the ranged accessor is multi-dimensional, the sub-range is allowed to +/// > describe a region of memory in the underlying buffer that is not +/// > contiguous in the linear address space. +/// > +/// > Most of the accessor member functions which provide a reference to the +/// > underlying buffer elements are affected by a ranged accessor’s offset and +/// > range. ... In addition, the accessor’s iterator functions iterate only +/// > over the elements that are within the sub-range. +/// +/// Classes below implement the logic of iterating through N-dimensional +/// (1 <= N <= 3) space, which covers a potentially non-contiguous memory +/// region in the underlying accessor bufffer. +/// +/// Most of the logic is implemented in __accessor_iterator_base class, which +/// provides routines for all the indexing logic such as +/// incrementing/decrementing iterators, addition/substraction and comparison +/// operators of iterators, etc. +/// +/// Pointer to accessor is held by __accessor_iterator class, which provides +/// user-visible interface of iterator. + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { +template +class accessor; + +namespace detail { + +/// Base class for accessor iterators, which implements common logic between +/// all iterators (const, reverse, const reverse, etc.) +/// +/// In order to iterate through a possibly non-contiguous N-dimensional space, +/// the class holds an N-dimensional `id`, which is carefuly incremented each +/// time iterator is incremented/decrementing, taking into account the +/// shape/size of a space iterator goes through. +/// +/// Whilst increment/decrement operation can be implemented through a couple of +/// 'if's and assignments, additon/substraction operators which can move an +/// iterator up to N elements, are harder to implement on a N-dimensional id. +/// In order to implement them, the class also holds and maintains a linearized +/// id, which can be quickly updated to perform an addition/substraction of an +/// iterator. However, that id has to be deleniarized in order to be used to +/// dereference particular element of an accessor and that operation includes +/// division and taking reminder of the division. Those operations are more +/// expensive than simple additional and conditionals and therefore the class +/// maintains both N-dimensional and linear id to balance between implementation +/// simplicity and performance of (presumably) most oftenly used operations with +/// an accessor. +template +class __accessor_iterator_base { +protected: + using difference_type = size_t; + using iterator_category = std::random_access_iterator_tag; + +private: + id<_Dimensions> _MBegin; + // Holds an id which is relative to _MBegin. + id<_Dimensions> _MCurrent; + id<_Dimensions> _MEnd; + + static constexpr int _Index0 = _Dimensions - 1; + static constexpr int _Index1 = _Dimensions - 2; + static constexpr int _Index2 = _Dimensions - 3; + + static constexpr difference_type _LinearBegin = 0; + // Holds an id which is relative to _LinearBegin + difference_type _MLinearCurrent = 0; + difference_type _MLinearEnd = 0; + + difference_type _MRowSize = 0; + difference_type _MSliceSize = 0; + +protected: + __accessor_iterator_base() {} + + __accessor_iterator_base(const id<_Dimensions> &_Begin, + const id<_Dimensions> &_End, + const id<_Dimensions> &_Current) + : _MBegin(_Begin), _MCurrent(_Current - _MBegin), _MEnd(_End) { + _MLinearEnd = _MRowSize = _MEnd[_Index0] - _MBegin[_Index0]; + if constexpr (_Dimensions > 1) { + _MSliceSize = (_MEnd[_Index1] - _MBegin[_Index1]) * _MRowSize; + // Multiply by number of rows + _MLinearEnd *= _MEnd[_Index1] - _MBegin[_Index1]; + } + if constexpr (_Dimensions > 2) { + // Multiply by number of slices + _MLinearEnd *= _MEnd[_Index2] - _MBegin[_Index2]; + } + _MLinearCurrent = __linearizeIndex(_MCurrent); + } + + id<_Dimensions> __get_current_id() const { + return _MBegin + _MCurrent; + } + + __accessor_iterator_base &operator++() { + if constexpr (_IsReverse) + __decrement(); + else + __increment(); + return *this; + } + + __accessor_iterator_base operator++(int) { + auto _Old = *this; + ++(*this); + return _Old; + } + + __accessor_iterator_base &operator--() { + if constexpr (_IsReverse) + __increment(); + else + __decrement(); + return *this; + } + + __accessor_iterator_base operator--(int) { + auto _Old = *this; + --(*this); + return _Old; + } + + __accessor_iterator_base &operator+=(difference_type _N) { + // iterator && N > 0 -> forward + // iterator && N < 0 -> backwards + // reverse iterator && N > 0 -> backwards + // reverse iterator && N < 0 -> forward + bool _BackwardsDirection = !_IsReverse ^ (_N > 0); + if (_BackwardsDirection) + __adjustBackwards(_N); + else + __adjustForward(_N); + return *this; + } + + __accessor_iterator_base &operator-=(difference_type _N) { + // iterator && N > 0 -> backwards + // iterator && N < 0 -> forward + // reverse iterator && N > 0 -> forward + // reverse iterator && N < 0 -> backwards + bool _ForwardDirection = !_IsReverse ^ (_N > 0); + if (_ForwardDirection) + __adjustForward(_N); + else + __adjustBackwards(_N); + return *this; + } + + difference_type operator-(const __accessor_iterator_base &_Rhs) { + if (_Rhs._MLinearCurrent > _MLinearCurrent) + return _Rhs._MLinearCurrent - _MLinearCurrent; + else + return _MLinearCurrent - _Rhs._MLinearCurrent; + } + + bool operator<(const __accessor_iterator_base<_Dimensions> &_Other) const { + return _MLinearCurrent < _Other._MLinearCurrent; + } + + bool operator>(const __accessor_iterator_base<_Dimensions> &_Other) const { + return _Other < *this; + } + + bool operator<=(const __accessor_iterator_base<_Dimensions> &_Other) const { + return !(*this > _Other); + } + + bool operator>=(const __accessor_iterator_base<_Dimensions> &_Other) const { + return !(*this < _Other); + } + + bool operator==(const __accessor_iterator_base<_Dimensions> &_Other) const { + return _MLinearCurrent == _Other._MLinearCurrent; + } + + bool operator!=(const __accessor_iterator_base<_Dimensions> &_Other) const { + return !(*this == _Other); + } + +private: + void __increment() { + if (_MLinearCurrent >= _MLinearEnd) + return; + + ++_MLinearCurrent; + if (_MCurrent[_Index0] < _MEnd[_Index0]) + _MCurrent[_Index0]++; + if constexpr (_Dimensions > 1) { + if (_MCurrent[_Index0] == _MEnd[_Index0]) { + if (_MCurrent[_Index1] < _MEnd[_Index1]) { + _MCurrent[_Index1]++; + _MCurrent[_Index0] = _MBegin[_Index0]; + } + } + } + if constexpr (_Dimensions > 2) { + if (_MCurrent[_Index1] == _MEnd[_Index1]) { + if (_MCurrent[_Index2] < _MEnd[_Index2]) { + _MCurrent[_Index2]++; + _MCurrent[_Index0] = _MBegin[_Index0]; + _MCurrent[_Index1] = _MBegin[_Index1]; + } + } + } + } + + void __decrement() { + if (_MLinearCurrent == _LinearBegin) + return; + + --_MLinearCurrent; + if (_MCurrent[_Index0] > 0) + _MCurrent[_Index0]--; + if constexpr (_Dimensions > 1) { + if (_MCurrent[_Index0] == 0) { + if (_MCurrent[_Index1] > 0) { + _MCurrent[_Index1]--; + _MCurrent[_Index0] = _MEnd[_Index0] - 1; + } + } + } + if constexpr (_Dimensions > 2) { + if (_MCurrent[_Index1] == 0) { + if (_MCurrent[_Index2] > 0) { + _MCurrent[_Index2]--; + _MCurrent[_Index0] = _MEnd[_Index0] - 1; + _MCurrent[_Index1] = _MEnd[_Index1] - 1; + } + } + } + } + + void __adjustForward(difference_type _N) { + if (_MLinearCurrent + _N > _MLinearEnd) + _MLinearCurrent = _MLinearEnd; + else + _MLinearCurrent += _N; + _MCurrent = __delinearizeIndex(_MLinearCurrent); + } + + void __adjustBackwards(difference_type _N) { + if (_N > _MLinearCurrent) + _MLinearCurrent = _LinearBegin; + else + _MLinearCurrent -= _N; + _MCurrent = __delinearizeIndex(_MLinearCurrent); + } + + size_t __linearizeIndex(const id<_Dimensions> &_Id) const { + size_t _Result = _Id[_Index0]; + if constexpr (_Dimensions > 1) + _Result += _Id[_Index1] * _MRowSize; + if constexpr (_Dimensions > 2) + _Result += _Id[_Index2] * _MSliceSize; + return _Result; + } + + id<_Dimensions> __delinearizeIndex(size_t _LinearId) const { + id<_Dimensions> _Result; + if constexpr (_Dimensions > 2) { + _Result[_Index2] = _LinearId / _MSliceSize; + _LinearId %= _MSliceSize; + } + if constexpr (_Dimensions > 1) { + _Result[_Index1] = _LinearId / _MRowSize; + _LinearId %= _MRowSize; + } + _Result[_Index0] = _LinearId; + return _Result; + } +}; + +template +class __accessor_iterator : public __accessor_iterator_base<_Dimensions> { + using _AccessorT = accessor<_DataT, _Dimensions, _AccessMode, _AccessTarget, + _IsPlaceholder, _PropertyListT>; + _AccessorT *_MAccessorPtr; + + using _BaseT = __accessor_iterator_base<_Dimensions>; + + friend class accessor<_DataT, _Dimensions, _AccessMode, _AccessTarget, + _IsPlaceholder, _PropertyListT>; + + __accessor_iterator(_AccessorT *_AccessorPtr, const id<_Dimensions> &_Begin, + const id<_Dimensions> &_End, + const id<_Dimensions> &_Current) + : __accessor_iterator_base<_Dimensions>(_Begin, _End, _Current), + _MAccessorPtr(_AccessorPtr) {} + + static __accessor_iterator __get_begin(_AccessorT *_AccessorPtr, + const id<_Dimensions> &_Begin, + const id<_Dimensions> &_End) { + return __accessor_iterator(_AccessorPtr, _Begin, _End, _Begin); + } + + static __accessor_iterator __get_end(_AccessorT *_AccessorPtr, + const id<_Dimensions> &_Begin, + const id<_Dimensions> &_End) { + // As `.end()` iterator we use an iterator which points to the first element + // past the end of an accessible range. That is done to simplify the process + // of transforming an iterator to an `.end()` state by incrementing it. + // + // However, `_End` id passed here highlights an accessible range and do not + // point to the first element past the end of the accessible range in all + // cases. For example, let's take a look at a case where we access a + // 2-dimensional buffer of size 2x2. Inputs to this method will be: + // _Begin: (0, 0; _End(2, 2): + // Begin Elem . + // Elem Elem . + // . . End + // + // As showed above, _End simply defines the shape/size, but it doesn't point + // to the element we would like it to point to. That happens because _End + // passed here comes from an accessor range, which is 1-indexed. However, + // accessor::operator[] accepts a 0-indexed id. In order to create a + // past-the-end iterator, we convert _End id to a 0-indexed one, + // create an interator out of it and then simply increment it. + auto _EndCopy = _End; + for (auto _I = 0; _I < _Dimensions; ++_I) + _EndCopy[_I]--; + + auto _Ret = __accessor_iterator(_AccessorPtr, _Begin, _End, _EndCopy); + return ++_Ret; + } + +public: + using difference_type = typename _BaseT::difference_type; + using value_type = _DataT; + // FIXME: this should likely include address space + using pointer = _DataT *; + using reference = _DataT &; + using iterator_category = typename _BaseT::iterator_category; + + __accessor_iterator() : _MAccessorPtr(nullptr) {} + + _DataT &operator*() { + return _MAccessorPtr->operator[](this->__get_current_id()); + } + + __accessor_iterator &operator++() { + _BaseT::operator++(); + return *this; + } + + __accessor_iterator operator++(int) { + auto _Old = *this; + _BaseT::operator++(); + return _Old; + } + + __accessor_iterator &operator--() { + _BaseT::operator--(); + return *this; + } + + __accessor_iterator operator--(int) { + auto _Old = *this; + _BaseT::operator--(); + return _Old; + } + + __accessor_iterator &operator+=(difference_type _N) { + _BaseT::operator+=(_N); + return *this; + } + + friend __accessor_iterator operator+(__accessor_iterator &_Lhs, + difference_type _N) { + _Lhs += _N; + return _Lhs; + } + + friend __accessor_iterator operator+(difference_type _N, + __accessor_iterator &_Rhs) { + _Rhs += _N; + return _Rhs; + } + + __accessor_iterator &operator-=(difference_type _N) { + _BaseT::operator-=(_N); + return *this; + } + + friend __accessor_iterator operator-(__accessor_iterator &_Lhs, + difference_type _N) { + _Lhs -= _N; + return _Lhs; + } + + reference &operator[](difference_type _N) { + auto _Copy = *this; + _Copy += _N; + return *_Copy; + } + + using __accessor_iterator_base<_Dimensions>::operator-; + using __accessor_iterator_base<_Dimensions>::operator==; + using __accessor_iterator_base<_Dimensions>::operator!=; + using __accessor_iterator_base<_Dimensions>::operator<; + using __accessor_iterator_base<_Dimensions>::operator<=; + using __accessor_iterator_base<_Dimensions>::operator>; + using __accessor_iterator_base<_Dimensions>::operator>=; +}; +} // namespace detail +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/test/basic_tests/accessor/iterator.cpp b/sycl/test/basic_tests/accessor/iterator.cpp new file mode 100644 index 0000000000000..431ef793b7009 --- /dev/null +++ b/sycl/test/basic_tests/accessor/iterator.cpp @@ -0,0 +1,35 @@ +// RUN: %clangxx -fsycl %s -o %t.out +// RUN: %t.out + +#include + +#include + +using namespace sycl; + +int main() { + std::vector reference = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18}; + + { + buffer buf(reference.data(), range<1>{10}); + auto acc = buf.get_access(range<1>{10}); + std::vector data; + auto It = acc.begin(); + std::cout << *(It--) << std::endl; + std::cout << *(It--) << std::endl; + std::cout << *(--It) << std::endl; + std::cout << *(--It) << std::endl; + /*int N = 0; + for (auto I = acc.begin(), E = acc.end(); I != E; ++I) { + data.push_back(*I); + std::cout << *I << std::endl; + ++N; + if (N > 20) + break; + }*/ + + // assert + } + + return 0; +} diff --git a/sycl/unittests/CMakeLists.txt b/sycl/unittests/CMakeLists.txt index 39b0921629336..34fba725d88c1 100644 --- a/sycl/unittests/CMakeLists.txt +++ b/sycl/unittests/CMakeLists.txt @@ -46,3 +46,4 @@ add_subdirectory(Extensions) add_subdirectory(windows) add_subdirectory(event) add_subdirectory(buffer) +add_subdirectory(accessor) diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp new file mode 100644 index 0000000000000..74b2af79f114f --- /dev/null +++ b/sycl/unittests/accessor/AccessorIterator.cpp @@ -0,0 +1,232 @@ +#include + +#include + +#include +#include +#include +#include + +class AccessorIteratorTest : public ::testing::Test { +public: + AccessorIteratorTest() {} + + template + void checkFullCopyThroughIterator(const sycl::range &shape) { + std::vector reference(shape.size()); + std::iota(reference.begin(), reference.end(), 0); + sycl::buffer buffer(reference.data(), shape); + auto accessor = buffer.template get_access(); + std::vector copied; + auto I = accessor.begin(); + I = accessor.end(); + for (auto i = accessor.begin(), e = accessor.end(); i != e; ++i) { + copied.push_back(*i); + } + + ASSERT_EQ(copied.size(), reference.size()); + for (size_t i = 0, e = reference.size(); i < e; ++i) { + ASSERT_EQ(copied[i], reference[i]); + } + } + + template + void checkPartialCopyThroughIteratorWithoutOffset( + const sycl::range &fullShape, + const sycl::range ©Shape) { + std::vector reference(fullShape.size()); + std::iota(reference.begin(), reference.end(), 0); + sycl::buffer buffer(reference.data(), fullShape); + std::vector copied; + { + auto accessor = + buffer.template get_access(copyShape); + auto I = accessor.begin(); + I = accessor.end(); + for (auto i = accessor.begin(), e = accessor.end(); i != e; ++i) { + copied.push_back(*i); + } + } + + ASSERT_EQ(copied.size(), copyShape.size()); + + { + auto fullAccessor = buffer.template get_access(); + + if constexpr (Dimensions == 1) { + for (size_t x = 0; x < copyShape[0]; ++x) { + ASSERT_EQ(copied[x], reference[x]); + } + } else if constexpr (Dimensions == 2) { + size_t linear = 0; + for (size_t y = 0; y < copyShape[0]; ++y) { + for (size_t x = 0; x < copyShape[1]; ++x) { + ASSERT_EQ(copied[linear], fullAccessor[y][x]); + ++linear; + } + } + } else { + size_t linear = 0; + for (size_t z = 0; z < copyShape[0]; ++z) { + for (size_t y = 0; y < copyShape[1]; ++y) { + for (size_t x = 0; x < copyShape[2]; ++x) { + ASSERT_EQ(copied[linear], fullAccessor[z][y][x]); + ++linear; + } + } + } + } + } + } +}; + +TEST_F(AccessorIteratorTest, ImplementationDetails) { + std::vector reference(5); + std::iota(reference.begin(), reference.end(), 0); + sycl::buffer buffer(reference.data(), sycl::range<1>{reference.size()}); + auto accessor = buffer.template get_access(); + { + auto It = accessor.begin(); + // Check that It can't be decremented past begin + ASSERT_EQ(--It, accessor.begin()); + ASSERT_EQ(It - 1, accessor.begin()); + ASSERT_EQ(It -= 1, accessor.begin()); + ASSERT_EQ(It - 10, accessor.begin()); + ASSERT_EQ(It -= 10, accessor.begin()); + } + { + auto It = accessor.end(); + // Check that It can't be incremented past end + ASSERT_EQ(++It, accessor.end()); + ASSERT_EQ(It + 1, accessor.end()); + ASSERT_EQ(It += 1, accessor.end()); + ASSERT_EQ(It + 10, accessor.end()); + ASSERT_EQ(It += 10, accessor.end()); + } +} + +// FIXME: consider turning this into parameterized test to check various +// accessor types +TEST_F(AccessorIteratorTest, IteratorTraits) { + using IteratorT = sycl::accessor::iterator; + ASSERT_TRUE( + (std::is_same_v::difference_type, + std::iterator_traits::difference_type>)); + ASSERT_TRUE((std::is_same_v::value_type, + std::iterator_traits::value_type>)); + ASSERT_TRUE((std::is_same_v::value_type *, + std::iterator_traits::pointer>)); + ASSERT_TRUE((std::is_same_v::reference, + std::iterator_traits::reference>)); + ASSERT_TRUE( + (std::is_same_v::iterator_category>)); +} + +// Based on requirements listed at +// https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator +TEST_F(AccessorIteratorTest, LegacyRandomAccessIteratorRequirements) { + using IteratorT = sycl::accessor::iterator; + IteratorT It; + auto &RefToIt = It; + ASSERT_TRUE((std::is_same_v)); + ASSERT_TRUE((std::is_same_v)); + ASSERT_TRUE((std::is_same_v)); + ASSERT_TRUE((std::is_same_v)); + ASSERT_TRUE((std::is_same_v)); + IteratorT It2; + ASSERT_TRUE((std::is_same_v::difference_type, + decltype(It - It2)>)); + ASSERT_TRUE( + (std::is_convertible_v::reference>)); + ASSERT_TRUE((std::is_convertible_v)); + ASSERT_TRUE((std::is_convertible_v)); + ASSERT_TRUE((std::is_convertible_v It2), bool>)); + ASSERT_TRUE((std::is_convertible_v= It2), bool>)); + // FIXME: add more test cases based on Notes +} + +// Based on requirements listed at +// https://en.cppreference.com/w/cpp/named_req/BidirectionalIterator +TEST_F(AccessorIteratorTest, LegacyBidirectionalIteratorRequirements) { + using IteratorT = sycl::accessor::iterator; + IteratorT It; + ASSERT_TRUE((std::is_same_v)); + ASSERT_TRUE((std::is_convertible_v)); + ASSERT_TRUE((std::is_same_v::reference, + decltype(*It--)>)); +} + +// Based on requirements listed at +// https://en.cppreference.com/w/cpp/named_req/ForwardIterator +TEST_F(AccessorIteratorTest, LegacyForwardIteratorRequirements) { + using IteratorT = sycl::accessor::iterator; + ASSERT_TRUE(std::is_default_constructible_v); + IteratorT It; + ASSERT_TRUE((std::is_same_v)); + ASSERT_TRUE((std::is_same_v::reference, + decltype(*It++)>)); + IteratorT It2; + ASSERT_TRUE((std::is_convertible_v)); + ASSERT_TRUE((std::is_convertible_v)); + // FIXME: test that Objects of the type IteratorT provide multipass guarantee +} + +// Based on requirements listead at +// https://en.cppreference.com/w/cpp/named_req/Iterator +TEST_F(AccessorIteratorTest, LegacyIteratorRequirements) { + using IteratorT = sycl::accessor::iterator; + ASSERT_TRUE(std::is_copy_constructible_v); + ASSERT_TRUE(std::is_copy_assignable_v); + ASSERT_TRUE(std::is_destructible_v); + ASSERT_TRUE(std::is_swappable_v); + IteratorT It; + ASSERT_TRUE((std::is_same_v)); + ASSERT_TRUE((std::is_same_v::reference, + decltype(*It)>)); +} + +TEST_F(AccessorIteratorTest, FullCopy1D) { + ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIterator(sycl::range<1>{10})); +} + +TEST_F(AccessorIteratorTest, FullCopy2D) { + ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIterator(sycl::range<2>{2, 5})); + ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIterator(sycl::range<2>{5, 2})); + ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIterator(sycl::range<2>{1, 10})); + ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIterator(sycl::range<2>{10, 1})); +} + +TEST_F(AccessorIteratorTest, FullCopy3D) { + ASSERT_NO_FATAL_FAILURE( + checkFullCopyThroughIterator(sycl::range<3>{3, 3, 3})); + ASSERT_NO_FATAL_FAILURE( + checkFullCopyThroughIterator(sycl::range<3>{1, 3, 3})); + ASSERT_NO_FATAL_FAILURE( + checkFullCopyThroughIterator(sycl::range<3>{3, 1, 3})); + ASSERT_NO_FATAL_FAILURE( + checkFullCopyThroughIterator(sycl::range<3>{3, 3, 1})); +} + +TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset1D) { + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<1>{10}, sycl::range<1>{5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<1>{10}, sycl::range<1>{10})); +} + +TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset2D) { + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<2>{5, 5}, sycl::range<2>{3, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<2>{5, 5}, sycl::range<2>{5, 5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<2>{5, 5}, sycl::range<2>{2, 5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<2>{5, 5}, sycl::range<2>{5, 2})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<2>{5, 5}, sycl::range<2>{3, 2})); +} + +TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset3D) { + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 3, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 5, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 5, 5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{1, 2, 3})); +} diff --git a/sycl/unittests/accessor/CMakeLists.txt b/sycl/unittests/accessor/CMakeLists.txt new file mode 100644 index 0000000000000..8080fd43cfe81 --- /dev/null +++ b/sycl/unittests/accessor/CMakeLists.txt @@ -0,0 +1,3 @@ +add_sycl_unittest(AccessorTests OBJECT + AccessorIterator.cpp +) From 2af129dfb20e6ae84d38fde41e31af64d653b3f6 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Mon, 19 Sep 2022 10:52:34 -0400 Subject: [PATCH 02/21] Add more tests and fix operator+ --- .../include/sycl/detail/accessor_iterator.hpp | 14 +- sycl/test/basic_tests/accessor/iterator.cpp | 10 +- sycl/unittests/accessor/AccessorIterator.cpp | 121 +++++++++++++++--- 3 files changed, 120 insertions(+), 25 deletions(-) diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index 1e868f0783e6f..8549504ded344 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -395,16 +395,18 @@ class __accessor_iterator : public __accessor_iterator_base<_Dimensions> { return *this; } - friend __accessor_iterator operator+(__accessor_iterator &_Lhs, + friend __accessor_iterator operator+(const __accessor_iterator &_Lhs, difference_type _N) { - _Lhs += _N; - return _Lhs; + auto _Ret = _Lhs; + _Ret += _N; + return _Ret; } friend __accessor_iterator operator+(difference_type _N, - __accessor_iterator &_Rhs) { - _Rhs += _N; - return _Rhs; + const __accessor_iterator &_Rhs) { + auto _Ret = _Rhs; + _Ret += _N; + return _Ret; } __accessor_iterator &operator-=(difference_type _N) { diff --git a/sycl/test/basic_tests/accessor/iterator.cpp b/sycl/test/basic_tests/accessor/iterator.cpp index 431ef793b7009..9329e6c4083c3 100644 --- a/sycl/test/basic_tests/accessor/iterator.cpp +++ b/sycl/test/basic_tests/accessor/iterator.cpp @@ -15,10 +15,12 @@ int main() { auto acc = buf.get_access(range<1>{10}); std::vector data; auto It = acc.begin(); - std::cout << *(It--) << std::endl; - std::cout << *(It--) << std::endl; - std::cout << *(--It) << std::endl; - std::cout << *(--It) << std::endl; + It += 3; + std::cout << "loop start" << std::endl; + for (int i = -3; i <=3; ++i) { + std::cout << "i = " << i << std::endl; + std::cout << ((It + i) == (i + It)) << std::endl; + } /*int N = 0; for (auto I = acc.begin(), E = acc.end(); I != E; ++I) { data.push_back(*I); diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp index 74b2af79f114f..732b5b991fa28 100644 --- a/sycl/unittests/accessor/AccessorIterator.cpp +++ b/sycl/unittests/accessor/AccessorIterator.cpp @@ -144,7 +144,67 @@ TEST_F(AccessorIteratorTest, LegacyRandomAccessIteratorRequirements) { ASSERT_TRUE((std::is_convertible_v)); ASSERT_TRUE((std::is_convertible_v It2), bool>)); ASSERT_TRUE((std::is_convertible_v= It2), bool>)); - // FIXME: add more test cases based on Notes +} + +// Based on notes listed at +// https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator +TEST_F(AccessorIteratorTest, LegacyRandomAccessIteratorRequirementsExtra) { + std::vector reference(6); + std::iota(reference.begin(), reference.end(), 0); + sycl::buffer buffer(reference.data(), sycl::range<1>{reference.size()}); + auto accessor = buffer.template get_access(); + auto It = accessor.begin(); + It += 3; + + { // It += n should be equivalent to incrementint/decrementing It n times + for (int n = -3; n <= 3; ++n) { + auto It1 = It; + auto It2 = It; + It1 += n; + + if (n < 0) { + int i = n; + while (i++) + --It2; + } else { + int i = n; + while (i--) + ++It2; + } + + ASSERT_EQ(It1, It2); + } + } + + { // It + n == n + It + for (int n = -3; n <= 3; ++n) { + ASSERT_EQ(It + n, n + It); + } + } + + { + auto It1 = accessor.begin(); + auto It2 = accessor.end(); + ASSERT_EQ(It - It1, It1 - It); + ASSERT_EQ(It - It2, It2 - It); + ASSERT_EQ(It2, It + (It2 - It)); + ASSERT_EQ(It, It1 + (It - It1)); + } + + { + auto It1 = accessor.begin(); + auto It2 = accessor.begin(); + auto It3 = accessor.end(); + + ASSERT_TRUE(!(It1 < It2)); + ASSERT_TRUE(It1 < It); // precondition for the next check + ASSERT_TRUE(!(It < It1)); + ASSERT_TRUE(It < It3); // precondition for the next check + ASSERT_TRUE(It1 < It3); + + ASSERT_FALSE(It3 < It); + ASSERT_FALSE(It == It3); + } } // Based on requirements listed at @@ -170,7 +230,25 @@ TEST_F(AccessorIteratorTest, LegacyForwardIteratorRequirements) { IteratorT It2; ASSERT_TRUE((std::is_convertible_v)); ASSERT_TRUE((std::is_convertible_v)); - // FIXME: test that Objects of the type IteratorT provide multipass guarantee +} + +TEST_F(AccessorIteratorTest, MultipassGuarantee) { + std::vector reference(5); + std::iota(reference.begin(), reference.end(), 0); + sycl::buffer buffer(reference.data(), sycl::range<1>{reference.size()}); + auto accessor = buffer.template get_access(); + auto It1 = accessor.begin(); + auto It2 = accessor.begin(); + + while (It1 != accessor.end()) { + ASSERT_EQ(It1, It2); + ASSERT_EQ(*It1, *It2); + ASSERT_EQ(++It1, ++It2); + } + + It1 = accessor.begin(); + It2 = It1; + ASSERT_EQ(((void)++It2, *It1), *It1); } // Based on requirements listead at @@ -210,23 +288,36 @@ TEST_F(AccessorIteratorTest, FullCopy3D) { } TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset1D) { - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<1>{10}, sycl::range<1>{5})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<1>{10}, sycl::range<1>{10})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<1>{10}, sycl::range<1>{5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<1>{10}, sycl::range<1>{10})); } TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset2D) { - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<2>{5, 5}, sycl::range<2>{3, 3})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<2>{5, 5}, sycl::range<2>{5, 5})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<2>{5, 5}, sycl::range<2>{2, 5})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<2>{5, 5}, sycl::range<2>{5, 2})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<2>{5, 5}, sycl::range<2>{3, 2})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<2>{5, 5}, sycl::range<2>{3, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<2>{5, 5}, sycl::range<2>{5, 5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<2>{5, 5}, sycl::range<2>{2, 5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<2>{5, 5}, sycl::range<2>{5, 2})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<2>{5, 5}, sycl::range<2>{3, 2})); } TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset3D) { - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 3})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 3, 3})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 5, 3})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 5})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 5, 5})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset(sycl::range<3>{5, 5, 5}, sycl::range<3>{1, 2, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 3, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 5, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 5, 5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + sycl::range<3>{5, 5, 5}, sycl::range<3>{1, 2, 3})); } From d68418552d633a57fc40aa720a1a9c4fc5bd7d0a Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 21 Sep 2022 08:37:49 -0400 Subject: [PATCH 03/21] Mark begin/end methods as const --- sycl/include/sycl/accessor.hpp | 4 ++-- sycl/include/sycl/detail/accessor_iterator.hpp | 13 ++++++------- 2 files changed, 8 insertions(+), 9 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 9ca08196198a6..421e3dfff52dd 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -2029,11 +2029,11 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; } bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); } - iterator begin() noexcept { + iterator begin() const noexcept { return iterator::__get_begin(this, get_offset(), get_range()); } - iterator end() noexcept { + iterator end() const noexcept { return iterator::__get_end(this, get_offset(), get_range()); } diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index 8549504ded344..1d06ae8eba00e 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -117,9 +117,7 @@ class __accessor_iterator_base { _MLinearCurrent = __linearizeIndex(_MCurrent); } - id<_Dimensions> __get_current_id() const { - return _MBegin + _MCurrent; - } + id<_Dimensions> __get_current_id() const { return _MBegin + _MCurrent; } __accessor_iterator_base &operator++() { if constexpr (_IsReverse) @@ -305,26 +303,27 @@ template { using _AccessorT = accessor<_DataT, _Dimensions, _AccessMode, _AccessTarget, _IsPlaceholder, _PropertyListT>; - _AccessorT *_MAccessorPtr; + const _AccessorT *_MAccessorPtr; using _BaseT = __accessor_iterator_base<_Dimensions>; friend class accessor<_DataT, _Dimensions, _AccessMode, _AccessTarget, _IsPlaceholder, _PropertyListT>; - __accessor_iterator(_AccessorT *_AccessorPtr, const id<_Dimensions> &_Begin, + __accessor_iterator(const _AccessorT *_AccessorPtr, + const id<_Dimensions> &_Begin, const id<_Dimensions> &_End, const id<_Dimensions> &_Current) : __accessor_iterator_base<_Dimensions>(_Begin, _End, _Current), _MAccessorPtr(_AccessorPtr) {} - static __accessor_iterator __get_begin(_AccessorT *_AccessorPtr, + static __accessor_iterator __get_begin(const _AccessorT *_AccessorPtr, const id<_Dimensions> &_Begin, const id<_Dimensions> &_End) { return __accessor_iterator(_AccessorPtr, _Begin, _End, _Begin); } - static __accessor_iterator __get_end(_AccessorT *_AccessorPtr, + static __accessor_iterator __get_end(const _AccessorT *_AccessorPtr, const id<_Dimensions> &_Begin, const id<_Dimensions> &_End) { // As `.end()` iterator we use an iterator which points to the first element From 2dff81dc73ee29119e26430f50fe5daaf02b1f6f Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 21 Sep 2022 13:05:40 -0400 Subject: [PATCH 04/21] Generalize test code a bit --- sycl/unittests/accessor/AccessorIterator.cpp | 47 +++++++++++--------- 1 file changed, 26 insertions(+), 21 deletions(-) diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp index 732b5b991fa28..cee4d1eb080f6 100644 --- a/sycl/unittests/accessor/AccessorIterator.cpp +++ b/sycl/unittests/accessor/AccessorIterator.cpp @@ -30,6 +30,22 @@ class AccessorIteratorTest : public ::testing::Test { } } + template + auto &&accessHelper(Container &&C, int Idx, Indices... Ids) { + if constexpr (CurrentDimension > TotalDimensions) { + (void)Idx; + return accessHelper(C, Ids...); + } else + return accessHelper(C[Idx], + Ids...); + } + + template + auto &&accessHelper(Container &&C, int Idx) { + return C[Idx]; + } + template void checkPartialCopyThroughIteratorWithoutOffset( const sycl::range &fullShape, @@ -52,27 +68,16 @@ class AccessorIteratorTest : public ::testing::Test { { auto fullAccessor = buffer.template get_access(); - - if constexpr (Dimensions == 1) { - for (size_t x = 0; x < copyShape[0]; ++x) { - ASSERT_EQ(copied[x], reference[x]); - } - } else if constexpr (Dimensions == 2) { - size_t linear = 0; - for (size_t y = 0; y < copyShape[0]; ++y) { - for (size_t x = 0; x < copyShape[1]; ++x) { - ASSERT_EQ(copied[linear], fullAccessor[y][x]); - ++linear; - } - } - } else { - size_t linear = 0; - for (size_t z = 0; z < copyShape[0]; ++z) { - for (size_t y = 0; y < copyShape[1]; ++y) { - for (size_t x = 0; x < copyShape[2]; ++x) { - ASSERT_EQ(copied[linear], fullAccessor[z][y][x]); - ++linear; - } + size_t linearId = 0; + sycl::id<3> shapeToCheck(Dimensions > 2 ? copyShape[Dimensions - 3] : 1, + Dimensions > 1 ? copyShape[Dimensions - 2] : 1, + copyShape[Dimensions - 1]); + for (size_t z = 0; z < shapeToCheck[0]; ++z) { + for (size_t y = 0; y < shapeToCheck[1]; ++y) { + for (size_t x = 0; x < shapeToCheck[2]; ++x) { + auto value = accessHelper(fullAccessor, z, y, x); + ASSERT_EQ(copied[linearId], value); + ++linearId; } } } From 87e771fbf707298533e8f1f4bf337ff993ad3f47 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 21 Sep 2022 14:18:54 -0400 Subject: [PATCH 05/21] Fixes and tests for offset + range accessors --- sycl/include/sycl/accessor.hpp | 5 +- sycl/unittests/accessor/AccessorIterator.cpp | 96 ++++++++++++++------ 2 files changed, 73 insertions(+), 28 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 421e3dfff52dd..8438eaa5f32d8 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -2030,11 +2030,12 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); } iterator begin() const noexcept { - return iterator::__get_begin(this, get_offset(), get_range()); + return iterator::__get_begin(this, get_offset(), + get_offset() + get_range()); } iterator end() const noexcept { - return iterator::__get_end(this, get_offset(), get_range()); + return iterator::__get_end(this, get_offset(), get_offset() + get_range()); } private: diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp index cee4d1eb080f6..c0ede121122ac 100644 --- a/sycl/unittests/accessor/AccessorIterator.cpp +++ b/sycl/unittests/accessor/AccessorIterator.cpp @@ -47,16 +47,17 @@ class AccessorIteratorTest : public ::testing::Test { } template - void checkPartialCopyThroughIteratorWithoutOffset( - const sycl::range &fullShape, - const sycl::range ©Shape) { + void + checkPartialCopyThroughIterator(const sycl::range &fullShape, + const sycl::range ©Shape, + const sycl::id &offset = {}) { std::vector reference(fullShape.size()); std::iota(reference.begin(), reference.end(), 0); sycl::buffer buffer(reference.data(), fullShape); std::vector copied; { - auto accessor = - buffer.template get_access(copyShape); + auto accessor = buffer.template get_access( + copyShape, offset); auto I = accessor.begin(); I = accessor.end(); for (auto i = accessor.begin(), e = accessor.end(); i != e; ++i) { @@ -69,12 +70,19 @@ class AccessorIteratorTest : public ::testing::Test { { auto fullAccessor = buffer.template get_access(); size_t linearId = 0; - sycl::id<3> shapeToCheck(Dimensions > 2 ? copyShape[Dimensions - 3] : 1, - Dimensions > 1 ? copyShape[Dimensions - 2] : 1, - copyShape[Dimensions - 1]); - for (size_t z = 0; z < shapeToCheck[0]; ++z) { - for (size_t y = 0; y < shapeToCheck[1]; ++y) { - for (size_t x = 0; x < shapeToCheck[2]; ++x) { + + sycl::id<3> offsetToUse(Dimensions > 2 ? offset[Dimensions - 3] : 1, + Dimensions > 1 ? offset[Dimensions - 2] : 1, + offset[Dimensions - 1]); + + sycl::id<3> shapeToCheck( + (Dimensions > 2 ? copyShape[Dimensions - 3] : 1) - offsetToUse[0], + (Dimensions > 1 ? copyShape[Dimensions - 2] : 1) - offsetToUse[1], + copyShape[Dimensions - 1] - offsetToUse[2]); + + for (size_t z = offsetToUse[0]; z < shapeToCheck[0]; ++z) { + for (size_t y = offsetToUse[1]; y < shapeToCheck[1]; ++y) { + for (size_t x = offsetToUse[2]; x < shapeToCheck[2]; ++x) { auto value = accessHelper(fullAccessor, z, y, x); ASSERT_EQ(copied[linearId], value); ++linearId; @@ -293,36 +301,72 @@ TEST_F(AccessorIteratorTest, FullCopy3D) { } TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset1D) { - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( - sycl::range<1>{10}, sycl::range<1>{5})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( - sycl::range<1>{10}, sycl::range<1>{10})); + ASSERT_NO_FATAL_FAILURE( + checkPartialCopyThroughIterator(sycl::range<1>{10}, sycl::range<1>{5})); + ASSERT_NO_FATAL_FAILURE( + checkPartialCopyThroughIterator(sycl::range<1>{10}, sycl::range<1>{10})); } TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset2D) { - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<2>{5, 5}, sycl::range<2>{3, 3})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<2>{5, 5}, sycl::range<2>{5, 5})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<2>{5, 5}, sycl::range<2>{2, 5})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<2>{5, 5}, sycl::range<2>{5, 2})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<2>{5, 5}, sycl::range<2>{3, 2})); } TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset3D) { - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 3})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 3, 3})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 5, 3})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 5})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 5, 5})); - ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIteratorWithoutOffset( + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<3>{5, 5, 5}, sycl::range<3>{1, 2, 3})); } + +TEST_F(AccessorIteratorTest, PartialCopyWithOffset1D) { + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<1>{10}, sycl::range<1>{5}, sycl::id<1>{3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<1>{10}, sycl::range<1>{5}, sycl::id<1>{5})); +} + +TEST_F(AccessorIteratorTest, PartialCopyWithOffset2D) { + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<2>{10, 10}, sycl::range<2>{5, 5}, sycl::id<2>{3, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<2>{10, 10}, sycl::range<2>{5, 5}, sycl::id<2>{3, 0})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<2>{10, 10}, sycl::range<2>{5, 5}, sycl::id<2>{0, 3})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<2>{10, 10}, sycl::range<2>{5, 5}, sycl::id<2>{5, 5})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<2>{5, 10}, sycl::range<2>{3, 5}, sycl::id<2>{1, 4})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<2>{10, 5}, sycl::range<2>{5, 3}, sycl::id<2>{5, 1})); +} + +TEST_F(AccessorIteratorTest, PartialCopyWithOffset3D) { + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<3>{7, 7, 7}, sycl::range<3>{3, 3, 3}, sycl::id<3>{2, 2, 2})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<3>{8, 8, 8}, sycl::range<3>{4, 4, 4}, sycl::id<3>{4, 4, 4})); + // FIXME: figure out why the test below fails + // ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + // sycl::range<3>{7, 7, 7}, sycl::range<3>{3, 3, 3}, sycl::id<3>{4, 4, 4})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<3>{7, 7, 7}, sycl::range<3>{3, 4, 5}, sycl::id<3>{3, 2, 1})); + ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( + sycl::range<3>{9, 8, 7}, sycl::range<3>{3, 4, 5}, sycl::id<3>{3, 2, 1})); +} From 36364a02817b9320dbc5c997e928c70a79876e03 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 22 Sep 2022 09:16:25 -0400 Subject: [PATCH 06/21] Remove LIT which was used for local experiments --- sycl/test/basic_tests/accessor/iterator.cpp | 37 --------------------- 1 file changed, 37 deletions(-) delete mode 100644 sycl/test/basic_tests/accessor/iterator.cpp diff --git a/sycl/test/basic_tests/accessor/iterator.cpp b/sycl/test/basic_tests/accessor/iterator.cpp deleted file mode 100644 index 9329e6c4083c3..0000000000000 --- a/sycl/test/basic_tests/accessor/iterator.cpp +++ /dev/null @@ -1,37 +0,0 @@ -// RUN: %clangxx -fsycl %s -o %t.out -// RUN: %t.out - -#include - -#include - -using namespace sycl; - -int main() { - std::vector reference = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18}; - - { - buffer buf(reference.data(), range<1>{10}); - auto acc = buf.get_access(range<1>{10}); - std::vector data; - auto It = acc.begin(); - It += 3; - std::cout << "loop start" << std::endl; - for (int i = -3; i <=3; ++i) { - std::cout << "i = " << i << std::endl; - std::cout << ((It + i) == (i + It)) << std::endl; - } - /*int N = 0; - for (auto I = acc.begin(), E = acc.end(); I != E; ++I) { - data.push_back(*I); - std::cout << *I << std::endl; - ++N; - if (N > 20) - break; - }*/ - - // assert - } - - return 0; -} From ba4557784b2781a6b750b3489e1d6cf850b4a615 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Mon, 26 Sep 2022 07:28:37 -0400 Subject: [PATCH 07/21] Completely refactor accessor::iterator Now iterators for non-ranged accessors should be almost equal to plain pointers and only iterator dereference for a ranged accessor is a costly operation which requires various calculations including division. --- sycl/include/sycl/accessor.hpp | 5 +- .../include/sycl/detail/accessor_iterator.hpp | 577 ++++++++---------- sycl/unittests/accessor/AccessorIterator.cpp | 11 +- 3 files changed, 264 insertions(+), 329 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 8438eaa5f32d8..4f3ad7d122148 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -2030,12 +2030,11 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); } iterator begin() const noexcept { - return iterator::__get_begin(this, get_offset(), - get_offset() + get_range()); + return iterator::__get_begin(this, getMemoryRange()); } iterator end() const noexcept { - return iterator::__get_end(this, get_offset(), get_offset() + get_range()); + return iterator::__get_end(this, getMemoryRange()); } private: diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index 1d06ae8eba00e..873e0fc5a327a 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -10,7 +10,9 @@ #include +#include #include +#include #include /// \file accessor_iterator.hpp @@ -31,18 +33,6 @@ /// > underlying buffer elements are affected by a ranged accessor’s offset and /// > range. ... In addition, the accessor’s iterator functions iterate only /// > over the elements that are within the sub-range. -/// -/// Classes below implement the logic of iterating through N-dimensional -/// (1 <= N <= 3) space, which covers a potentially non-contiguous memory -/// region in the underlying accessor bufffer. -/// -/// Most of the logic is implemented in __accessor_iterator_base class, which -/// provides routines for all the indexing logic such as -/// incrementing/decrementing iterators, addition/substraction and comparison -/// operators of iterators, etc. -/// -/// Pointer to accessor is held by __accessor_iterator class, which provides -/// user-visible interface of iterator. namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { @@ -53,385 +43,330 @@ class accessor; namespace detail { -/// Base class for accessor iterators, which implements common logic between -/// all iterators (const, reverse, const reverse, etc.) -/// -/// In order to iterate through a possibly non-contiguous N-dimensional space, -/// the class holds an N-dimensional `id`, which is carefuly incremented each -/// time iterator is incremented/decrementing, taking into account the -/// shape/size of a space iterator goes through. -/// -/// Whilst increment/decrement operation can be implemented through a couple of -/// 'if's and assignments, additon/substraction operators which can move an -/// iterator up to N elements, are harder to implement on a N-dimensional id. -/// In order to implement them, the class also holds and maintains a linearized -/// id, which can be quickly updated to perform an addition/substraction of an -/// iterator. However, that id has to be deleniarized in order to be used to -/// dereference particular element of an accessor and that operation includes -/// division and taking reminder of the division. Those operations are more -/// expensive than simple additional and conditionals and therefore the class -/// maintains both N-dimensional and linear id to balance between implementation -/// simplicity and performance of (presumably) most oftenly used operations with -/// an accessor. -template -class __accessor_iterator_base { -protected: - using difference_type = size_t; +template +class __accessor_iterator { +public: + using difference_type = std::ptrdiff_t; + using value_type = _DataT; + // FIXME: this should likely include address space + using pointer = _DataT *; + using reference = _DataT &; using iterator_category = std::random_access_iterator_tag; -private: - id<_Dimensions> _MBegin; - // Holds an id which is relative to _MBegin. - id<_Dimensions> _MCurrent; - id<_Dimensions> _MEnd; - - static constexpr int _Index0 = _Dimensions - 1; - static constexpr int _Index1 = _Dimensions - 2; - static constexpr int _Index2 = _Dimensions - 3; - - static constexpr difference_type _LinearBegin = 0; - // Holds an id which is relative to _LinearBegin - difference_type _MLinearCurrent = 0; - difference_type _MLinearEnd = 0; - - difference_type _MRowSize = 0; - difference_type _MSliceSize = 0; - -protected: - __accessor_iterator_base() {} - - __accessor_iterator_base(const id<_Dimensions> &_Begin, - const id<_Dimensions> &_End, - const id<_Dimensions> &_Current) - : _MBegin(_Begin), _MCurrent(_Current - _MBegin), _MEnd(_End) { - _MLinearEnd = _MRowSize = _MEnd[_Index0] - _MBegin[_Index0]; - if constexpr (_Dimensions > 1) { - _MSliceSize = (_MEnd[_Index1] - _MBegin[_Index1]) * _MRowSize; - // Multiply by number of rows - _MLinearEnd *= _MEnd[_Index1] - _MBegin[_Index1]; - } - if constexpr (_Dimensions > 2) { - // Multiply by number of slices - _MLinearEnd *= _MEnd[_Index2] - _MBegin[_Index2]; - } - _MLinearCurrent = __linearizeIndex(_MCurrent); - } + __accessor_iterator() = default; - id<_Dimensions> __get_current_id() const { return _MBegin + _MCurrent; } + _DataT &operator*() { + return *(_MAccessorPtr->get_pointer() + __get_absolute_offset_to_buffer()); + } - __accessor_iterator_base &operator++() { - if constexpr (_IsReverse) - __decrement(); - else - __increment(); + __accessor_iterator &operator++() { + if (_MLinearId < _MEnd) + ++_MLinearId; return *this; } - __accessor_iterator_base operator++(int) { + __accessor_iterator operator++(int) { auto _Old = *this; ++(*this); return _Old; } - __accessor_iterator_base &operator--() { - if constexpr (_IsReverse) - __increment(); - else - __decrement(); + __accessor_iterator &operator--() { + if (_MLinearId > _MBegin) + --_MLinearId; return *this; } - __accessor_iterator_base operator--(int) { + __accessor_iterator operator--(int) { auto _Old = *this; --(*this); return _Old; } - __accessor_iterator_base &operator+=(difference_type _N) { - // iterator && N > 0 -> forward - // iterator && N < 0 -> backwards - // reverse iterator && N > 0 -> backwards - // reverse iterator && N < 0 -> forward - bool _BackwardsDirection = !_IsReverse ^ (_N > 0); - if (_BackwardsDirection) - __adjustBackwards(_N); - else - __adjustForward(_N); - return *this; - } + __accessor_iterator &operator+=(difference_type _N) { + if (_N < 0) { + *this -= -_N; + return *this; + } - __accessor_iterator_base &operator-=(difference_type _N) { - // iterator && N > 0 -> backwards - // iterator && N < 0 -> forward - // reverse iterator && N > 0 -> forward - // reverse iterator && N < 0 -> backwards - bool _ForwardDirection = !_IsReverse ^ (_N > 0); - if (_ForwardDirection) - __adjustForward(_N); + if (static_cast(_N) > _MEnd || _MEnd - _N < _MLinearId) + _MLinearId = _MEnd; else - __adjustBackwards(_N); + _MLinearId += _N; + return *this; } - difference_type operator-(const __accessor_iterator_base &_Rhs) { - if (_Rhs._MLinearCurrent > _MLinearCurrent) - return _Rhs._MLinearCurrent - _MLinearCurrent; - else - return _MLinearCurrent - _Rhs._MLinearCurrent; + friend __accessor_iterator operator+(const __accessor_iterator &_Lhs, + difference_type _N) { + auto _Ret = _Lhs; + _Ret += _N; + return _Ret; } - bool operator<(const __accessor_iterator_base<_Dimensions> &_Other) const { - return _MLinearCurrent < _Other._MLinearCurrent; + friend __accessor_iterator operator+(difference_type _N, + const __accessor_iterator &_Rhs) { + auto _Ret = _Rhs; + _Ret += _N; + return _Ret; } - bool operator>(const __accessor_iterator_base<_Dimensions> &_Other) const { - return _Other < *this; - } + __accessor_iterator &operator-=(difference_type _N) { + if (_N < 0) { + *this += -_N; + return *this; + } - bool operator<=(const __accessor_iterator_base<_Dimensions> &_Other) const { - return !(*this > _Other); + if (_MBegin + _N > _MLinearId) + _MLinearId = _MBegin; + else + _MLinearId -= _N; + + return *this; } - bool operator>=(const __accessor_iterator_base<_Dimensions> &_Other) const { - return !(*this < _Other); + friend __accessor_iterator operator-(__accessor_iterator &_Lhs, + difference_type _N) { + _Lhs -= _N; + return _Lhs; } - bool operator==(const __accessor_iterator_base<_Dimensions> &_Other) const { - return _MLinearCurrent == _Other._MLinearCurrent; + reference &operator[](difference_type _N) { + auto _Copy = *this; + _Copy += _N; + return *_Copy; } - bool operator!=(const __accessor_iterator_base<_Dimensions> &_Other) const { - return !(*this == _Other); + bool operator<(const __accessor_iterator &_Other) const { + return _MLinearId < _Other._MLinearId; } -private: - void __increment() { - if (_MLinearCurrent >= _MLinearEnd) - return; - - ++_MLinearCurrent; - if (_MCurrent[_Index0] < _MEnd[_Index0]) - _MCurrent[_Index0]++; - if constexpr (_Dimensions > 1) { - if (_MCurrent[_Index0] == _MEnd[_Index0]) { - if (_MCurrent[_Index1] < _MEnd[_Index1]) { - _MCurrent[_Index1]++; - _MCurrent[_Index0] = _MBegin[_Index0]; - } - } - } - if constexpr (_Dimensions > 2) { - if (_MCurrent[_Index1] == _MEnd[_Index1]) { - if (_MCurrent[_Index2] < _MEnd[_Index2]) { - _MCurrent[_Index2]++; - _MCurrent[_Index0] = _MBegin[_Index0]; - _MCurrent[_Index1] = _MBegin[_Index1]; - } - } - } + bool operator>(const __accessor_iterator &_Other) const { + return _Other < *this; } - void __decrement() { - if (_MLinearCurrent == _LinearBegin) - return; - - --_MLinearCurrent; - if (_MCurrent[_Index0] > 0) - _MCurrent[_Index0]--; - if constexpr (_Dimensions > 1) { - if (_MCurrent[_Index0] == 0) { - if (_MCurrent[_Index1] > 0) { - _MCurrent[_Index1]--; - _MCurrent[_Index0] = _MEnd[_Index0] - 1; - } - } - } - if constexpr (_Dimensions > 2) { - if (_MCurrent[_Index1] == 0) { - if (_MCurrent[_Index2] > 0) { - _MCurrent[_Index2]--; - _MCurrent[_Index0] = _MEnd[_Index0] - 1; - _MCurrent[_Index1] = _MEnd[_Index1] - 1; - } - } - } + bool operator<=(const __accessor_iterator &_Other) const { + return !(*this > _Other); } - void __adjustForward(difference_type _N) { - if (_MLinearCurrent + _N > _MLinearEnd) - _MLinearCurrent = _MLinearEnd; - else - _MLinearCurrent += _N; - _MCurrent = __delinearizeIndex(_MLinearCurrent); + bool operator>=(const __accessor_iterator &_Other) const { + return !(*this < _Other); } - void __adjustBackwards(difference_type _N) { - if (_N > _MLinearCurrent) - _MLinearCurrent = _LinearBegin; - else - _MLinearCurrent -= _N; - _MCurrent = __delinearizeIndex(_MLinearCurrent); + bool operator==(const __accessor_iterator &_Other) const { + return _MLinearId == _Other._MLinearId; } - size_t __linearizeIndex(const id<_Dimensions> &_Id) const { - size_t _Result = _Id[_Index0]; - if constexpr (_Dimensions > 1) - _Result += _Id[_Index1] * _MRowSize; - if constexpr (_Dimensions > 2) - _Result += _Id[_Index2] * _MSliceSize; - return _Result; + bool operator!=(const __accessor_iterator &_Other) const { + return !(*this == _Other); } - id<_Dimensions> __delinearizeIndex(size_t _LinearId) const { - id<_Dimensions> _Result; - if constexpr (_Dimensions > 2) { - _Result[_Index2] = _LinearId / _MSliceSize; - _LinearId %= _MSliceSize; - } - if constexpr (_Dimensions > 1) { - _Result[_Index1] = _LinearId / _MRowSize; - _LinearId %= _MRowSize; - } - _Result[_Index0] = _LinearId; - return _Result; + difference_type operator-(const __accessor_iterator &_Rhs) { + // FIXME: values of difference_type can be negative + if (_Rhs._MLinearId > _MLinearId) + return _Rhs._MLinearId - _MLinearId; + else + return _MLinearId - _Rhs._MLinearId; } -}; -template -class __accessor_iterator : public __accessor_iterator_base<_Dimensions> { +private: using _AccessorT = accessor<_DataT, _Dimensions, _AccessMode, _AccessTarget, _IsPlaceholder, _PropertyListT>; - const _AccessorT *_MAccessorPtr; - - using _BaseT = __accessor_iterator_base<_Dimensions>; - friend class accessor<_DataT, _Dimensions, _AccessMode, _AccessTarget, _IsPlaceholder, _PropertyListT>; - __accessor_iterator(const _AccessorT *_AccessorPtr, - const id<_Dimensions> &_Begin, - const id<_Dimensions> &_End, - const id<_Dimensions> &_Current) - : __accessor_iterator_base<_Dimensions>(_Begin, _End, _Current), - _MAccessorPtr(_AccessorPtr) {} - - static __accessor_iterator __get_begin(const _AccessorT *_AccessorPtr, - const id<_Dimensions> &_Begin, - const id<_Dimensions> &_End) { - return __accessor_iterator(_AccessorPtr, _Begin, _End, _Begin); - } - - static __accessor_iterator __get_end(const _AccessorT *_AccessorPtr, - const id<_Dimensions> &_Begin, - const id<_Dimensions> &_End) { - // As `.end()` iterator we use an iterator which points to the first element - // past the end of an accessible range. That is done to simplify the process - // of transforming an iterator to an `.end()` state by incrementing it. - // - // However, `_End` id passed here highlights an accessible range and do not - // point to the first element past the end of the accessible range in all - // cases. For example, let's take a look at a case where we access a - // 2-dimensional buffer of size 2x2. Inputs to this method will be: - // _Begin: (0, 0; _End(2, 2): - // Begin Elem . - // Elem Elem . - // . . End - // - // As showed above, _End simply defines the shape/size, but it doesn't point - // to the element we would like it to point to. That happens because _End - // passed here comes from an accessor range, which is 1-indexed. However, - // accessor::operator[] accepts a 0-indexed id. In order to create a - // past-the-end iterator, we convert _End id to a 0-indexed one, - // create an interator out of it and then simply increment it. - auto _EndCopy = _End; - for (auto _I = 0; _I < _Dimensions; ++_I) - _EndCopy[_I]--; - - auto _Ret = __accessor_iterator(_AccessorPtr, _Begin, _End, _EndCopy); - return ++_Ret; - } - -public: - using difference_type = typename _BaseT::difference_type; - using value_type = _DataT; - // FIXME: this should likely include address space - using pointer = _DataT *; - using reference = _DataT &; - using iterator_category = typename _BaseT::iterator_category; - - __accessor_iterator() : _MAccessorPtr(nullptr) {} + const _AccessorT *_MAccessorPtr = nullptr; + + // Stores a linear id of an accessor's buffer element the iterator points to. + // This id is relative to a range accessible through an accessor, i.e. it is + // limited by a space with top left corner defiend as accessor::get_offset() + // and bottom right corner defined as accesor::get_range(). + size_t _MLinearId = 0; + + // Describes range of linear IDs accessible by the iterator. _MEnd corresponds + // to ID of en element past the last accessible element of accessors's + // buffer. + size_t _MBegin = 0; + size_t _MEnd = 0; + + // If set to true, then it indicates that accessor has its offset and/or range + // set to non-zero, i.e. it is a ranged accessor. + bool _MAccessorIsRanged = false; + + // Fields below are used (and changed to be non-zero) only if we deal with + // a ranged accessor. + // + // TODO: consider making their existance dependable on _Dimensions template + // parameter, because not all of them are needed for all possible dimensions. + + // Three field below allow us to calculate an absolute offset to an accessor's + // buffer to correctly identify a memory region which this iterator should + // point to. Comments below describe them using an iterator to the following + // accessor as an example: + // + // buffer buf(input.data(), range<2>{5, 5}); + // auto acc = buf.get_access(range<2>{3, 3}, id<2>{1, 1}); + // + // Such combination of buffer size, access range and offset is visualized + // below. Dot (.) symbols represent buffer elements NOT reacheable by the + // accessor; X symbols represent buffer elements which ARE reachable by the + // the accessor. + // + // . . . . . + // . X X X . + // . X X X . + // . X X X . + // . . . . . + // + // _MStaticOffset stores a number of elements in _full_ rows (and in _full_ + // slices in case of 3-dimensional buffers) before the first accessible + // element. For the example above, _MStaticOffset would be equal to 5, because + // there is only one full row before the first accessible element. "Static" in + // the name highlights that this is a constant element in an equation which + // calculates an absoulte offset to an accessor's buffer, it doesn't depend + // on the current state of the iterator. + // + // _MPerRowOffset stores a number of _inaccessible_ elements in each + // _accessible_ row. For the example above it would be equal to 2 (leftmost + // and the rightmost elements of a row). + // + // _MPerSliceOffset stores a number of _inaccessible_ elements in each + // _accessible_ slice. Slice here means a single 2D layer in a 3D buffer. For + // the example above it would be equal to 0, because we are not looking at a + // 3D buffer. However, if we had two slices like visualized above, + // _MPerSliceOffset would be equal to 16 (elements on the "perimeter" of the + // slice, i.e. ones represented as dots (.)). + + size_t _MStaticOffset = 0; + size_t _MPerRowOffset = 0; + size_t _MPerSliceOffset = 0; + + // Contains a number of _accessible_ elements in a row + size_t _MRowSize = 0; + // Contains a number of _accessible_ elements in a slice + size_t _MSliceSize = 0; + + // Contains a full range of the underlying buffer + range<3> _MAccessRange = range<3>{0, 0, 0}; + + // _MLinearId stores an offset which is relative to the accessible range of + // the accessor, which means that it could be the case that _MlinearId equal + // to 0 should not correspond to the beginning of the underlying buffer, but + // instead should be re-adjusted to account for an offset passed to the + // accessor constructor. + // + // This function performs necessary calculations to make sure that all + // access ranges and offsets are taken into account. + size_t __get_absolute_offset_to_buffer() { + // For 1D case, any possible offsets are already incorporated into + // _MLinearId, so 1D is always treated as a non-ranged accessor + if (!_MAccessorIsRanged || _Dimensions == 1) + return _MLinearId; + + // Here we need to deal with 2D or 3D ranged accessor. + // _MLinearId points to an element relative to the accessible range. It + // should be adjusted to account for elements which are outside of the + // accessible range of the accessor. + + // We start with static offset: that is a number of elements in full rows + // and full slices before the first accessible element. + size_t _AbsoluteId = _MLinearId + _MStaticOffset; + + // Then we account for inaccessible elements in each full slice + size_t _Remaining = _MLinearId; + if constexpr (_Dimensions == 3) { + _AbsoluteId += _MPerSliceOffset * (_Remaining / _MSliceSize); + _Remaining %= _MSliceSize; + } - _DataT &operator*() { - return _MAccessorPtr->operator[](this->__get_current_id()); - } + // Then we account for inaccessible elements in each full row + _AbsoluteId += _MPerRowOffset * (_Remaining / _MRowSize); + _Remaining %= _MRowSize; - __accessor_iterator &operator++() { - _BaseT::operator++(); - return *this; - } + // And finally, there could be inaccessible elements on the current row + _AbsoluteId += _MAccessorPtr->get_offset()[_Dimensions - 1]; - __accessor_iterator operator++(int) { - auto _Old = *this; - _BaseT::operator++(); - return _Old; + return _AbsoluteId; } - __accessor_iterator &operator--() { - _BaseT::operator--(); - return *this; - } - - __accessor_iterator operator--(int) { - auto _Old = *this; - _BaseT::operator--(); - return _Old; - } + __accessor_iterator(const _AccessorT *_AccessorPtr, + const range<3> &_AccessRange) + : _MAccessorPtr(_AccessorPtr), _MAccessRange(_AccessRange) { + constexpr int _XIndex = _Dimensions - 1; + constexpr int _YIndex = _Dimensions - 2; + (void)_YIndex; + constexpr int _ZIndex = _Dimensions - 3; + (void)_ZIndex; - __accessor_iterator &operator+=(difference_type _N) { - _BaseT::operator+=(_N); - return *this; - } + if constexpr (_Dimensions > 1) + _MRowSize = _MAccessorPtr->get_range()[_XIndex]; + if constexpr (_Dimensions > 2) + _MSliceSize = _MAccessorPtr->get_range()[_YIndex] * _MRowSize; + + if (id<_Dimensions>{} != _MAccessorPtr->get_offset()) + _MAccessorIsRanged = true; + else { + for (size_t _I = 0; _I < _Dimensions; ++_I) + if (_MAccessorPtr->get_range()[_I] != _MAccessRange[_I]) + _MAccessorIsRanged = true; + } - friend __accessor_iterator operator+(const __accessor_iterator &_Lhs, - difference_type _N) { - auto _Ret = _Lhs; - _Ret += _N; - return _Ret; - } + if (_MAccessorIsRanged) { + if constexpr (_Dimensions > 2) { + _MStaticOffset += _MAccessRange[_XIndex] * _MAccessRange[_YIndex] * + _MAccessorPtr->get_offset()[_ZIndex]; + _MPerSliceOffset = + _MAccessRange[_XIndex] * _MAccessRange[_YIndex] - _MSliceSize; + } + if constexpr (_Dimensions > 1) { + _MStaticOffset += + _MAccessRange[_XIndex] * _MAccessorPtr->get_offset()[_YIndex]; + _MPerRowOffset = _MAccessRange[_XIndex] - _MRowSize; + } + } - friend __accessor_iterator operator+(difference_type _N, - const __accessor_iterator &_Rhs) { - auto _Ret = _Rhs; - _Ret += _N; - return _Ret; - } + // To further optimize 1D case, offset is already included into _Begin + if constexpr (_Dimensions == 1) + _MBegin = _MAccessorPtr->get_offset()[_XIndex]; - __accessor_iterator &operator-=(difference_type _N) { - _BaseT::operator-=(_N); - return *this; + _MEnd = _MBegin + _MAccessorPtr->size(); } - friend __accessor_iterator operator-(__accessor_iterator &_Lhs, - difference_type _N) { - _Lhs -= _N; - return _Lhs; + static __accessor_iterator __get_begin(const _AccessorT *_AccessorPtr, + const range<3> &_AccessRange) { + auto _It = __accessor_iterator(_AccessorPtr, _AccessRange); + _It._MLinearId = _It._MBegin; + return _It; } - reference &operator[](difference_type _N) { - auto _Copy = *this; - _Copy += _N; - return *_Copy; + static __accessor_iterator __get_end(const _AccessorT *_AccessorPtr, + const range<3> &_AccessRange) { + auto _It = __accessor_iterator(_AccessorPtr, _AccessRange); + _It._MLinearId = _It._MEnd; + return _It; } - using __accessor_iterator_base<_Dimensions>::operator-; - using __accessor_iterator_base<_Dimensions>::operator==; - using __accessor_iterator_base<_Dimensions>::operator!=; - using __accessor_iterator_base<_Dimensions>::operator<; - using __accessor_iterator_base<_Dimensions>::operator<=; - using __accessor_iterator_base<_Dimensions>::operator>; - using __accessor_iterator_base<_Dimensions>::operator>=; +public: +#ifndef NDEBUG + // Could be useful for debugging, but not a part of the official API, + // therefore only available in builds with assertions enabled. + friend std::ostream &operator<<(std::ostream &os, + const __accessor_iterator &it) { + os << "__accessor_iterator {\n"; + os << "\t_MLinearId: " << it._MLinearId << "\n"; + os << "\t_MEnd: " << it._MEnd << "\n"; + os << "\t_MStaticOffset: " << it._MStaticOffset << "\n"; + os << "\t_MPerRowOffset: " << it._MPerRowOffset << "\n"; + os << "\t_MPerSliceOffset: " << it._MPerSliceOffset << "\n"; + os << "\t_MRowSize: " << it._MRowSize << "\n"; + os << "\t_MSliceSize: " << it._MSliceSize << "\n"; + os << "\t_MAccessorIsRanged: " << it._MAccessorIsRanged << "\n"; + os << "}"; + return os; + } +#endif // NDEBUG }; } // namespace detail } // __SYCL_INLINE_VER_NAMESPACE(_V1) diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp index c0ede121122ac..c0d44392dc84a 100644 --- a/sycl/unittests/accessor/AccessorIterator.cpp +++ b/sycl/unittests/accessor/AccessorIterator.cpp @@ -76,9 +76,9 @@ class AccessorIteratorTest : public ::testing::Test { offset[Dimensions - 1]); sycl::id<3> shapeToCheck( - (Dimensions > 2 ? copyShape[Dimensions - 3] : 1) - offsetToUse[0], - (Dimensions > 1 ? copyShape[Dimensions - 2] : 1) - offsetToUse[1], - copyShape[Dimensions - 1] - offsetToUse[2]); + (Dimensions > 2 ? copyShape[Dimensions - 3] : 1) + offsetToUse[0], + (Dimensions > 1 ? copyShape[Dimensions - 2] : 1) + offsetToUse[1], + copyShape[Dimensions - 1] + offsetToUse[2]); for (size_t z = offsetToUse[0]; z < shapeToCheck[0]; ++z) { for (size_t y = offsetToUse[1]; y < shapeToCheck[1]; ++y) { @@ -185,7 +185,7 @@ TEST_F(AccessorIteratorTest, LegacyRandomAccessIteratorRequirementsExtra) { ++It2; } - ASSERT_EQ(It1, It2); + ASSERT_EQ(It1, It2) << " with n = " << n; } } @@ -364,7 +364,8 @@ TEST_F(AccessorIteratorTest, PartialCopyWithOffset3D) { sycl::range<3>{8, 8, 8}, sycl::range<3>{4, 4, 4}, sycl::id<3>{4, 4, 4})); // FIXME: figure out why the test below fails // ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( - // sycl::range<3>{7, 7, 7}, sycl::range<3>{3, 3, 3}, sycl::id<3>{4, 4, 4})); + // sycl::range<3>{7, 7, 7}, sycl::range<3>{3, 3, 3}, sycl::id<3>{4, 4, + // 4})); ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<3>{7, 7, 7}, sycl::range<3>{3, 4, 5}, sycl::id<3>{3, 2, 1})); ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( From 53d5799382b99fe26a417d9d7b7ad063c77a2346 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 28 Sep 2022 09:15:17 -0400 Subject: [PATCH 08/21] Some fixes and a couple more tests --- sycl/include/sycl/accessor.hpp | 2 +- sycl/include/sycl/detail/accessor_iterator.hpp | 6 +----- sycl/unittests/accessor/AccessorIterator.cpp | 8 ++++++-- 3 files changed, 8 insertions(+), 8 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 4f3ad7d122148..22c42cf376007 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -31,8 +31,8 @@ #include #include -#include #include +#include #include diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index 873e0fc5a327a..9d733fd07a5bf 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -164,11 +164,7 @@ class __accessor_iterator { } difference_type operator-(const __accessor_iterator &_Rhs) { - // FIXME: values of difference_type can be negative - if (_Rhs._MLinearId > _MLinearId) - return _Rhs._MLinearId - _MLinearId; - else - return _MLinearId - _Rhs._MLinearId; + return _MLinearId - _Rhs._MLinearId; } private: diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp index c0d44392dc84a..3a7ddf3ee4c3f 100644 --- a/sycl/unittests/accessor/AccessorIterator.cpp +++ b/sycl/unittests/accessor/AccessorIterator.cpp @@ -170,6 +170,7 @@ TEST_F(AccessorIteratorTest, LegacyRandomAccessIteratorRequirementsExtra) { It += 3; { // It += n should be equivalent to incrementint/decrementing It n times + // The test also checks the same for operator +, i.e. It + n for (int n = -3; n <= 3; ++n) { auto It1 = It; auto It2 = It; @@ -186,6 +187,7 @@ TEST_F(AccessorIteratorTest, LegacyRandomAccessIteratorRequirementsExtra) { } ASSERT_EQ(It1, It2) << " with n = " << n; + ASSERT_EQ(It + n, It2) << " with n = " << n; } } @@ -198,8 +200,10 @@ TEST_F(AccessorIteratorTest, LegacyRandomAccessIteratorRequirementsExtra) { { auto It1 = accessor.begin(); auto It2 = accessor.end(); - ASSERT_EQ(It - It1, It1 - It); - ASSERT_EQ(It - It2, It2 - It); + ASSERT_EQ(std::abs(It - It1), std::abs(It1 - It)); + ASSERT_EQ(std::abs(It - It2), std::abs(It2 - It)); + ASSERT_EQ(It1 - It, -3); + ASSERT_EQ(It - It1, 3); ASSERT_EQ(It2, It + (It2 - It)); ASSERT_EQ(It, It1 + (It - It1)); } From 90004cd5bf50b7cabd8a407322577c89f2cf2668 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 28 Sep 2022 09:33:25 -0400 Subject: [PATCH 09/21] Refactor tests to prepare for their generalization --- sycl/unittests/accessor/AccessorIterator.cpp | 58 ++++++++++---------- 1 file changed, 29 insertions(+), 29 deletions(-) diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp index 3a7ddf3ee4c3f..946b44fad9314 100644 --- a/sycl/unittests/accessor/AccessorIterator.cpp +++ b/sycl/unittests/accessor/AccessorIterator.cpp @@ -9,20 +9,14 @@ class AccessorIteratorTest : public ::testing::Test { public: - AccessorIteratorTest() {} - template void checkFullCopyThroughIterator(const sycl::range &shape) { std::vector reference(shape.size()); std::iota(reference.begin(), reference.end(), 0); sycl::buffer buffer(reference.data(), shape); auto accessor = buffer.template get_access(); - std::vector copied; - auto I = accessor.begin(); - I = accessor.end(); - for (auto i = accessor.begin(), e = accessor.end(); i != e; ++i) { - copied.push_back(*i); - } + std::vector copied = + copyThroughIterators(accessor.begin(), accessor.end()); ASSERT_EQ(copied.size(), reference.size()); for (size_t i = 0, e = reference.size(); i < e; ++i) { @@ -30,22 +24,6 @@ class AccessorIteratorTest : public ::testing::Test { } } - template - auto &&accessHelper(Container &&C, int Idx, Indices... Ids) { - if constexpr (CurrentDimension > TotalDimensions) { - (void)Idx; - return accessHelper(C, Ids...); - } else - return accessHelper(C[Idx], - Ids...); - } - - template - auto &&accessHelper(Container &&C, int Idx) { - return C[Idx]; - } - template void checkPartialCopyThroughIterator(const sycl::range &fullShape, @@ -58,11 +36,7 @@ class AccessorIteratorTest : public ::testing::Test { { auto accessor = buffer.template get_access( copyShape, offset); - auto I = accessor.begin(); - I = accessor.end(); - for (auto i = accessor.begin(), e = accessor.end(); i != e; ++i) { - copied.push_back(*i); - } + copied = copyThroughIterators(accessor.begin(), accessor.end()); } ASSERT_EQ(copied.size(), copyShape.size()); @@ -91,6 +65,32 @@ class AccessorIteratorTest : public ::testing::Test { } } } + +private: + template + std::vector copyThroughIterators(IteratorT begin, IteratorT end) { + std::vector copied; + for (auto it = begin; it != end; ++it) { + copied.push_back(*it); + } + return copied; + } + + template + auto &&accessHelper(Container &&C, int Idx, Indices... Ids) { + if constexpr (CurrentDimension > TotalDimensions) { + (void)Idx; + return accessHelper(C, Ids...); + } else + return accessHelper(C[Idx], + Ids...); + } + + template + auto &&accessHelper(Container &&C, int Idx) { + return C[Idx]; + } }; TEST_F(AccessorIteratorTest, ImplementationDetails) { From 7ff5db5f091b914c243311ca95d8d3ea3ec033d5 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 28 Sep 2022 13:46:25 -0400 Subject: [PATCH 10/21] Do not store accessor memory range in iterator --- sycl/include/sycl/detail/accessor_iterator.hpp | 15 ++++++--------- 1 file changed, 6 insertions(+), 9 deletions(-) diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index 9d733fd07a5bf..1e2def7197a51 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -244,9 +244,6 @@ class __accessor_iterator { // Contains a number of _accessible_ elements in a slice size_t _MSliceSize = 0; - // Contains a full range of the underlying buffer - range<3> _MAccessRange = range<3>{0, 0, 0}; - // _MLinearId stores an offset which is relative to the accessible range of // the accessor, which means that it could be the case that _MlinearId equal // to 0 should not correspond to the beginning of the underlying buffer, but @@ -289,7 +286,7 @@ class __accessor_iterator { __accessor_iterator(const _AccessorT *_AccessorPtr, const range<3> &_AccessRange) - : _MAccessorPtr(_AccessorPtr), _MAccessRange(_AccessRange) { + : _MAccessorPtr(_AccessorPtr) { constexpr int _XIndex = _Dimensions - 1; constexpr int _YIndex = _Dimensions - 2; (void)_YIndex; @@ -305,21 +302,21 @@ class __accessor_iterator { _MAccessorIsRanged = true; else { for (size_t _I = 0; _I < _Dimensions; ++_I) - if (_MAccessorPtr->get_range()[_I] != _MAccessRange[_I]) + if (_MAccessorPtr->get_range()[_I] != _AccessRange[_I]) _MAccessorIsRanged = true; } if (_MAccessorIsRanged) { if constexpr (_Dimensions > 2) { - _MStaticOffset += _MAccessRange[_XIndex] * _MAccessRange[_YIndex] * + _MStaticOffset += _AccessRange[_XIndex] * _AccessRange[_YIndex] * _MAccessorPtr->get_offset()[_ZIndex]; _MPerSliceOffset = - _MAccessRange[_XIndex] * _MAccessRange[_YIndex] - _MSliceSize; + _AccessRange[_XIndex] * _AccessRange[_YIndex] - _MSliceSize; } if constexpr (_Dimensions > 1) { _MStaticOffset += - _MAccessRange[_XIndex] * _MAccessorPtr->get_offset()[_YIndex]; - _MPerRowOffset = _MAccessRange[_XIndex] - _MRowSize; + _AccessRange[_XIndex] * _MAccessorPtr->get_offset()[_YIndex]; + _MPerRowOffset = _AccessRange[_XIndex] - _MRowSize; } } From d785db42c9cc873ee1f991c0fa5b112c44486d70 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 28 Sep 2022 13:56:25 -0400 Subject: [PATCH 11/21] Remove uses of accessor::get_offset() from methods other than iterator constructor --- .../include/sycl/detail/accessor_iterator.hpp | 24 +++++++++++-------- 1 file changed, 14 insertions(+), 10 deletions(-) diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index 1e2def7197a51..5c61b28ba49b4 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -216,13 +216,17 @@ class __accessor_iterator { // . X X X . // . . . . . // - // _MStaticOffset stores a number of elements in _full_ rows (and in _full_ - // slices in case of 3-dimensional buffers) before the first accessible - // element. For the example above, _MStaticOffset would be equal to 5, because - // there is only one full row before the first accessible element. "Static" in - // the name highlights that this is a constant element in an equation which - // calculates an absoulte offset to an accessor's buffer, it doesn't depend - // on the current state of the iterator. + // _MStaticOffset stores a number of elements which precede the first + // accessible element, calculated as if the buffer was linearized. + // For the example above, _MStaticOffset would be equal to 6, because + // there is one full row before the first accessible element and a one more on + // the second line. "Static" in the name highlights that this is a constant + // element in an equation which calculates an absoulte offset to an accessor's + // buffer, it doesn't depend on the current state of the iterator. + // + // NOTE: _MStaticOffset is set to 0 in 1D case even if the accessor was + // created with offset: it is done to further optimize 1D case by + // incorporating that offset into _MLinearId right away. // // _MPerRowOffset stores a number of _inaccessible_ elements in each // _accessible_ row. For the example above it would be equal to 2 (leftmost @@ -278,9 +282,6 @@ class __accessor_iterator { _AbsoluteId += _MPerRowOffset * (_Remaining / _MRowSize); _Remaining %= _MRowSize; - // And finally, there could be inaccessible elements on the current row - _AbsoluteId += _MAccessorPtr->get_offset()[_Dimensions - 1]; - return _AbsoluteId; } @@ -314,8 +315,11 @@ class __accessor_iterator { _AccessRange[_XIndex] * _AccessRange[_YIndex] - _MSliceSize; } if constexpr (_Dimensions > 1) { + // Elements in fully inaccessible rows _MStaticOffset += _AccessRange[_XIndex] * _MAccessorPtr->get_offset()[_YIndex]; + // Elements from the first accessible row + _MStaticOffset += _MAccessorPtr->get_offset()[_XIndex]; _MPerRowOffset = _AccessRange[_XIndex] - _MRowSize; } } From 7e5789b555199a543960830779213a09219a03ea Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 28 Sep 2022 14:40:24 -0400 Subject: [PATCH 12/21] Do not use accessor APIs in the iterator constructor This is a prepartion commit for when the iterator won't hold a pointer to an accessor at all. --- sycl/include/sycl/accessor.hpp | 8 ++- .../include/sycl/detail/accessor_iterator.hpp | 52 +++++++++++-------- 2 files changed, 36 insertions(+), 24 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 22c42cf376007..401b4a4f2b3f2 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -2030,11 +2030,15 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); } iterator begin() const noexcept { - return iterator::__get_begin(this, getMemoryRange()); + return iterator::__get_begin( + this, detail::convertToArrayOfN(getMemoryRange()), + get_range(), get_offset()); } iterator end() const noexcept { - return iterator::__get_end(this, getMemoryRange()); + return iterator::__get_end( + this, detail::convertToArrayOfN(getMemoryRange()), + get_range(), get_offset()); } private: diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index 5c61b28ba49b4..b080819ca666a 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -286,7 +286,9 @@ class __accessor_iterator { } __accessor_iterator(const _AccessorT *_AccessorPtr, - const range<3> &_AccessRange) + const range<_Dimensions> &_MemoryRange, + const range<_Dimensions> &_AccessRange, + const id<_Dimensions> &_Offset) : _MAccessorPtr(_AccessorPtr) { constexpr int _XIndex = _Dimensions - 1; constexpr int _YIndex = _Dimensions - 2; @@ -295,52 +297,58 @@ class __accessor_iterator { (void)_ZIndex; if constexpr (_Dimensions > 1) - _MRowSize = _MAccessorPtr->get_range()[_XIndex]; + _MRowSize = _AccessRange[_XIndex]; if constexpr (_Dimensions > 2) - _MSliceSize = _MAccessorPtr->get_range()[_YIndex] * _MRowSize; + _MSliceSize = _AccessRange[_YIndex] * _MRowSize; - if (id<_Dimensions>{} != _MAccessorPtr->get_offset()) + if (id<_Dimensions>{} != _Offset) _MAccessorIsRanged = true; else { for (size_t _I = 0; _I < _Dimensions; ++_I) - if (_MAccessorPtr->get_range()[_I] != _AccessRange[_I]) + if (_AccessRange[_I] != _MemoryRange[_I]) _MAccessorIsRanged = true; } if (_MAccessorIsRanged) { if constexpr (_Dimensions > 2) { - _MStaticOffset += _AccessRange[_XIndex] * _AccessRange[_YIndex] * - _MAccessorPtr->get_offset()[_ZIndex]; + _MStaticOffset += + _MemoryRange[_XIndex] * _MemoryRange[_YIndex] * _Offset[_ZIndex]; _MPerSliceOffset = - _AccessRange[_XIndex] * _AccessRange[_YIndex] - _MSliceSize; + _MemoryRange[_XIndex] * _MemoryRange[_YIndex] - _MSliceSize; } if constexpr (_Dimensions > 1) { // Elements in fully inaccessible rows - _MStaticOffset += - _AccessRange[_XIndex] * _MAccessorPtr->get_offset()[_YIndex]; - // Elements from the first accessible row - _MStaticOffset += _MAccessorPtr->get_offset()[_XIndex]; - _MPerRowOffset = _AccessRange[_XIndex] - _MRowSize; + _MStaticOffset += _MemoryRange[_XIndex] * _Offset[_YIndex]; + _MPerRowOffset = _MemoryRange[_XIndex] - _MRowSize; } - } - // To further optimize 1D case, offset is already included into _Begin - if constexpr (_Dimensions == 1) - _MBegin = _MAccessorPtr->get_offset()[_XIndex]; + // Elements from the first accessible row + if constexpr (_Dimensions == 1) + // To further optimize 1D case, offset is already included into _Begin + _MBegin = _Offset[_XIndex]; + else + _MStaticOffset += _Offset[_XIndex]; + } - _MEnd = _MBegin + _MAccessorPtr->size(); + _MEnd = _MBegin + _AccessRange.size(); } static __accessor_iterator __get_begin(const _AccessorT *_AccessorPtr, - const range<3> &_AccessRange) { - auto _It = __accessor_iterator(_AccessorPtr, _AccessRange); + const range<_Dimensions> &_MemoryRange, + const range<_Dimensions> &_AccessRange, + const id<_Dimensions> &_Offset) { + auto _It = + __accessor_iterator(_AccessorPtr, _MemoryRange, _AccessRange, _Offset); _It._MLinearId = _It._MBegin; return _It; } static __accessor_iterator __get_end(const _AccessorT *_AccessorPtr, - const range<3> &_AccessRange) { - auto _It = __accessor_iterator(_AccessorPtr, _AccessRange); + const range<_Dimensions> &_MemoryRange, + const range<_Dimensions> &_AccessRange, + const id<_Dimensions> &_Offset) { + auto _It = + __accessor_iterator(_AccessorPtr, _MemoryRange, _AccessRange, _Offset); _It._MLinearId = _It._MEnd; return _It; } From a4a181705eda3726e11e81f50a6b1936c37a7249 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 28 Sep 2022 15:01:25 -0400 Subject: [PATCH 13/21] Do not store a reference to an accessor within iterator --- sycl/include/sycl/accessor.hpp | 15 ++++--- .../include/sycl/detail/accessor_iterator.hpp | 39 +++++++++---------- 2 files changed, 25 insertions(+), 29 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 401b4a4f2b3f2..3be5ac3e533c8 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -1151,10 +1151,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : using reference = DataT &; using const_reference = const DataT &; - using iterator = - typename detail::__accessor_iterator; + using iterator = typename detail::__accessor_iterator; using difference_type = typename std::iterator_traits::difference_type; @@ -2031,14 +2028,16 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : iterator begin() const noexcept { return iterator::__get_begin( - this, detail::convertToArrayOfN(getMemoryRange()), - get_range(), get_offset()); + get_pointer(), + detail::convertToArrayOfN(getMemoryRange()), get_range(), + get_offset()); } iterator end() const noexcept { return iterator::__get_end( - this, detail::convertToArrayOfN(getMemoryRange()), - get_range(), get_offset()); + get_pointer(), + detail::convertToArrayOfN(getMemoryRange()), get_range(), + get_offset()); } private: diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index b080819ca666a..9593227427b21 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -36,17 +36,15 @@ namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { -template + +template class accessor; namespace detail { -template -class __accessor_iterator { +template class __accessor_iterator { public: using difference_type = std::ptrdiff_t; using value_type = _DataT; @@ -57,8 +55,8 @@ class __accessor_iterator { __accessor_iterator() = default; - _DataT &operator*() { - return *(_MAccessorPtr->get_pointer() + __get_absolute_offset_to_buffer()); + reference operator*() { + return *(_MDataPtr + __get_absolute_offset_to_buffer()); } __accessor_iterator &operator++() { @@ -168,12 +166,12 @@ class __accessor_iterator { } private: - using _AccessorT = accessor<_DataT, _Dimensions, _AccessMode, _AccessTarget, - _IsPlaceholder, _PropertyListT>; - friend class accessor<_DataT, _Dimensions, _AccessMode, _AccessTarget, - _IsPlaceholder, _PropertyListT>; + template + friend class sycl::accessor; - const _AccessorT *_MAccessorPtr = nullptr; + _DataT *_MDataPtr = nullptr; // Stores a linear id of an accessor's buffer element the iterator points to. // This id is relative to a range accessible through an accessor, i.e. it is @@ -285,11 +283,10 @@ class __accessor_iterator { return _AbsoluteId; } - __accessor_iterator(const _AccessorT *_AccessorPtr, - const range<_Dimensions> &_MemoryRange, + __accessor_iterator(_DataT *_DataPtr, const range<_Dimensions> &_MemoryRange, const range<_Dimensions> &_AccessRange, const id<_Dimensions> &_Offset) - : _MAccessorPtr(_AccessorPtr) { + : _MDataPtr(_DataPtr) { constexpr int _XIndex = _Dimensions - 1; constexpr int _YIndex = _Dimensions - 2; (void)_YIndex; @@ -333,22 +330,22 @@ class __accessor_iterator { _MEnd = _MBegin + _AccessRange.size(); } - static __accessor_iterator __get_begin(const _AccessorT *_AccessorPtr, + static __accessor_iterator __get_begin(_DataT *_DataPtr, const range<_Dimensions> &_MemoryRange, const range<_Dimensions> &_AccessRange, const id<_Dimensions> &_Offset) { auto _It = - __accessor_iterator(_AccessorPtr, _MemoryRange, _AccessRange, _Offset); + __accessor_iterator(_DataPtr, _MemoryRange, _AccessRange, _Offset); _It._MLinearId = _It._MBegin; return _It; } - static __accessor_iterator __get_end(const _AccessorT *_AccessorPtr, + static __accessor_iterator __get_end(_DataT *_DataPtr, const range<_Dimensions> &_MemoryRange, const range<_Dimensions> &_AccessRange, const id<_Dimensions> &_Offset) { auto _It = - __accessor_iterator(_AccessorPtr, _MemoryRange, _AccessRange, _Offset); + __accessor_iterator(_DataPtr, _MemoryRange, _AccessRange, _Offset); _It._MLinearId = _It._MEnd; return _It; } From 47e6c46e39f5db7ff5b8b4e6e12a140fb7914bc9 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 28 Sep 2022 16:17:57 -0400 Subject: [PATCH 14/21] Add cbegin/cend and some tests for those methods --- sycl/include/sycl/accessor.hpp | 16 ++++ sycl/unittests/accessor/AccessorIterator.cpp | 85 +++++++++++++------- 2 files changed, 70 insertions(+), 31 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 3be5ac3e533c8..934b39aecb6be 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -1152,6 +1152,8 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : using const_reference = const DataT &; using iterator = typename detail::__accessor_iterator; + using const_iterator = + typename detail::__accessor_iterator; using difference_type = typename std::iterator_traits::difference_type; @@ -2040,6 +2042,20 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : get_offset()); } + const_iterator cbegin() const noexcept { + return const_iterator::__get_begin( + get_pointer(), + detail::convertToArrayOfN(getMemoryRange()), get_range(), + get_offset()); + } + + const_iterator cend() const noexcept { + return const_iterator::__get_end( + get_pointer(), + detail::convertToArrayOfN(getMemoryRange()), get_range(), + get_offset()); + } + private: #ifdef __SYCL_DEVICE_ONLY__ size_t getTotalOffset() const { diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp index 946b44fad9314..364d00c0a7b76 100644 --- a/sycl/unittests/accessor/AccessorIterator.cpp +++ b/sycl/unittests/accessor/AccessorIterator.cpp @@ -15,13 +15,11 @@ class AccessorIteratorTest : public ::testing::Test { std::iota(reference.begin(), reference.end(), 0); sycl::buffer buffer(reference.data(), shape); auto accessor = buffer.template get_access(); - std::vector copied = - copyThroughIterators(accessor.begin(), accessor.end()); - ASSERT_EQ(copied.size(), reference.size()); - for (size_t i = 0, e = reference.size(); i < e; ++i) { - ASSERT_EQ(copied[i], reference[i]); - } + ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIteratorImpl( + reference, accessor.begin(), accessor.end())); + ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIteratorImpl( + reference, accessor.cbegin(), accessor.cend())); } template @@ -33,49 +31,74 @@ class AccessorIteratorTest : public ::testing::Test { std::iota(reference.begin(), reference.end(), 0); sycl::buffer buffer(reference.data(), fullShape); std::vector copied; + { auto accessor = buffer.template get_access( copyShape, offset); copied = copyThroughIterators(accessor.begin(), accessor.end()); } - - ASSERT_EQ(copied.size(), copyShape.size()); + ASSERT_NO_FATAL_FAILURE( + validatePartialCopyThroughIterator(copied, buffer, copyShape, offset)); { - auto fullAccessor = buffer.template get_access(); - size_t linearId = 0; - - sycl::id<3> offsetToUse(Dimensions > 2 ? offset[Dimensions - 3] : 1, - Dimensions > 1 ? offset[Dimensions - 2] : 1, - offset[Dimensions - 1]); - - sycl::id<3> shapeToCheck( - (Dimensions > 2 ? copyShape[Dimensions - 3] : 1) + offsetToUse[0], - (Dimensions > 1 ? copyShape[Dimensions - 2] : 1) + offsetToUse[1], - copyShape[Dimensions - 1] + offsetToUse[2]); - - for (size_t z = offsetToUse[0]; z < shapeToCheck[0]; ++z) { - for (size_t y = offsetToUse[1]; y < shapeToCheck[1]; ++y) { - for (size_t x = offsetToUse[2]; x < shapeToCheck[2]; ++x) { - auto value = accessHelper(fullAccessor, z, y, x); - ASSERT_EQ(copied[linearId], value); - ++linearId; - } - } - } + auto accessor = buffer.template get_access( + copyShape, offset); + copied = copyThroughIterators(accessor.cbegin(), accessor.cend()); } + ASSERT_NO_FATAL_FAILURE( + validatePartialCopyThroughIterator(copied, buffer, copyShape, offset)); } private: + template + void checkFullCopyThroughIteratorImpl(const std::vector &reference, + IteratorT begin, IteratorT end) { + std::vector copied = copyThroughIterators(begin, end); + + ASSERT_EQ(copied.size(), reference.size()); + for (size_t i = 0, e = reference.size(); i < e; ++i) { + ASSERT_EQ(copied[i], reference[i]); + } + } + template std::vector copyThroughIterators(IteratorT begin, IteratorT end) { std::vector copied; - for (auto it = begin; it != end; ++it) { + for (auto it = begin; it != end; ++it) copied.push_back(*it); - } + return copied; } + template + void + validatePartialCopyThroughIterator(const std::vector &copied, + sycl::buffer &buffer, + const sycl::range ©Shape, + const sycl::id &offset = {}) { + auto fullAccessor = buffer.template get_access(); + size_t linearId = 0; + + sycl::id<3> offsetToUse(Dimensions > 2 ? offset[Dimensions - 3] : 1, + Dimensions > 1 ? offset[Dimensions - 2] : 1, + offset[Dimensions - 1]); + + sycl::id<3> shapeToCheck( + (Dimensions > 2 ? copyShape[Dimensions - 3] : 1) + offsetToUse[0], + (Dimensions > 1 ? copyShape[Dimensions - 2] : 1) + offsetToUse[1], + copyShape[Dimensions - 1] + offsetToUse[2]); + + for (size_t z = offsetToUse[0]; z < shapeToCheck[0]; ++z) { + for (size_t y = offsetToUse[1]; y < shapeToCheck[1]; ++y) { + for (size_t x = offsetToUse[2]; x < shapeToCheck[2]; ++x) { + auto value = accessHelper(fullAccessor, z, y, x); + ASSERT_EQ(copied[linearId], value); + ++linearId; + } + } + } + } + template auto &&accessHelper(Container &&C, int Idx, Indices... Ids) { From 9ef84c26753f60db1bebae4baefc067e29942765 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 4 Oct 2022 06:17:22 -0400 Subject: [PATCH 15/21] Fix build after merge --- sycl/include/sycl/accessor.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 381474eb1a45b..c663a6a3a5f48 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -1194,7 +1194,6 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : using value_type = DataT; using reference = DataT &; using const_reference = const DataT &; - using difference_type = size_t; using iterator = typename detail::__accessor_iterator; using const_iterator = From 91f0d30f9fda656f605564f4f4475d995a18ebcb Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Tue, 4 Oct 2022 08:35:51 -0400 Subject: [PATCH 16/21] Fix Wreorder warning coming from AccessorSubscript constructor --- sycl/include/sycl/accessor.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index c663a6a3a5f48..2b418b5acbd43 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -336,7 +336,7 @@ class accessor_common { public: AccessorSubscript(AccType Accessor, id IDs) - : MAccessor(Accessor), MIDs(IDs) {} + : MIDs(IDs), MAccessor(Accessor) {} // Only accessor class is supposed to use this c'tor for the first // operator[]. From 8e2db4ef3f3a6c95c161a5b4d0825ce1a3173fa0 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 7 Oct 2022 08:34:43 -0400 Subject: [PATCH 17/21] Allow iterators to be decremented past begin and incremented past end --- .../include/sycl/detail/accessor_iterator.hpp | 26 +++---------------- sycl/unittests/accessor/AccessorIterator.cpp | 25 ------------------ 2 files changed, 4 insertions(+), 47 deletions(-) diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index 9593227427b21..39b2f55ac846c 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -60,8 +60,7 @@ template class __accessor_iterator { } __accessor_iterator &operator++() { - if (_MLinearId < _MEnd) - ++_MLinearId; + ++_MLinearId; return *this; } @@ -72,8 +71,7 @@ template class __accessor_iterator { } __accessor_iterator &operator--() { - if (_MLinearId > _MBegin) - --_MLinearId; + --_MLinearId; return *this; } @@ -84,15 +82,7 @@ template class __accessor_iterator { } __accessor_iterator &operator+=(difference_type _N) { - if (_N < 0) { - *this -= -_N; - return *this; - } - - if (static_cast(_N) > _MEnd || _MEnd - _N < _MLinearId) - _MLinearId = _MEnd; - else - _MLinearId += _N; + _MLinearId += _N; return *this; } @@ -112,15 +102,7 @@ template class __accessor_iterator { } __accessor_iterator &operator-=(difference_type _N) { - if (_N < 0) { - *this += -_N; - return *this; - } - - if (_MBegin + _N > _MLinearId) - _MLinearId = _MBegin; - else - _MLinearId -= _N; + _MLinearId -= _N; return *this; } diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp index 364d00c0a7b76..0f2f0ed7fb039 100644 --- a/sycl/unittests/accessor/AccessorIterator.cpp +++ b/sycl/unittests/accessor/AccessorIterator.cpp @@ -116,31 +116,6 @@ class AccessorIteratorTest : public ::testing::Test { } }; -TEST_F(AccessorIteratorTest, ImplementationDetails) { - std::vector reference(5); - std::iota(reference.begin(), reference.end(), 0); - sycl::buffer buffer(reference.data(), sycl::range<1>{reference.size()}); - auto accessor = buffer.template get_access(); - { - auto It = accessor.begin(); - // Check that It can't be decremented past begin - ASSERT_EQ(--It, accessor.begin()); - ASSERT_EQ(It - 1, accessor.begin()); - ASSERT_EQ(It -= 1, accessor.begin()); - ASSERT_EQ(It - 10, accessor.begin()); - ASSERT_EQ(It -= 10, accessor.begin()); - } - { - auto It = accessor.end(); - // Check that It can't be incremented past end - ASSERT_EQ(++It, accessor.end()); - ASSERT_EQ(It + 1, accessor.end()); - ASSERT_EQ(It += 1, accessor.end()); - ASSERT_EQ(It + 10, accessor.end()); - ASSERT_EQ(It += 10, accessor.end()); - } -} - // FIXME: consider turning this into parameterized test to check various // accessor types TEST_F(AccessorIteratorTest, IteratorTraits) { From 769ab1380c8500dadd7e804a4836e4f6de1b8301 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 7 Oct 2022 09:35:30 -0400 Subject: [PATCH 18/21] Add tests for writing through an iterator --- sycl/unittests/accessor/AccessorIterator.cpp | 138 +++++++++++++++++++ 1 file changed, 138 insertions(+) diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp index 0f2f0ed7fb039..c67dbc7af51cd 100644 --- a/sycl/unittests/accessor/AccessorIterator.cpp +++ b/sycl/unittests/accessor/AccessorIterator.cpp @@ -9,6 +9,63 @@ class AccessorIteratorTest : public ::testing::Test { public: + template + void checkWriteThroughIterator(const sycl::range &fullShape, + const sycl::range &fillShape, + const sycl::id &offset) { + std::vector data(fullShape.size(), T{}); + sycl::buffer buffer(data.data(), fullShape); + { + auto accessor = buffer.template get_access( + fillShape, offset); + T linear_id = 1; + for (auto it = accessor.begin(), e = accessor.end(); it != e; ++it) { + *it = linear_id; + linear_id += 1; + } + } + + sycl::id<3> offsetToUse(Dimensions > 2 ? offset[Dimensions - 3] : 0, + Dimensions > 1 ? offset[Dimensions - 2] : 0, + offset[Dimensions - 1]); + + sycl::id<3> shapeToCheck( + (Dimensions > 2 ? fillShape[Dimensions - 3] : 1) + offsetToUse[0], + (Dimensions > 1 ? fillShape[Dimensions - 2] : 1) + offsetToUse[1], + fillShape[Dimensions - 1] + offsetToUse[2]); + + auto fullAccessor = buffer.template get_access(); + T linear_id = 1; + for (size_t z = offsetToUse[0]; z < shapeToCheck[0]; ++z) { + for (size_t y = offsetToUse[1]; y < shapeToCheck[1]; ++y) { + for (size_t x = offsetToUse[2]; x < shapeToCheck[2]; ++x) { + auto value = accessHelper(fullAccessor, z, y, x); + ASSERT_EQ(linear_id, value); + linear_id += 1; + } + } + } + + sycl::id<3> adjustedFullShape( + Dimensions > 2 ? fullShape[Dimensions - 3] : 1, + Dimensions > 1 ? fullShape[Dimensions - 2] : 1, + fullShape[Dimensions - 1]); + + for (size_t z = 0; z < adjustedFullShape[0]; ++z) { + for (size_t y = 0; y < adjustedFullShape[1]; ++y) { + for (size_t x = 0; x < adjustedFullShape[2]; ++x) { + // Skip elements which we previously checked + if (z >= offsetToUse[0] && z < shapeToCheck[0] && + y >= offsetToUse[1] && y < shapeToCheck[1] && + x >= offsetToUse[2] && x < shapeToCheck[2]) + continue; + auto value = accessHelper(fullAccessor, z, y, x); + ASSERT_EQ(T{}, value) << "at (" << z << "; " << y << "; " << x << ")"; + } + } + } + } + template void checkFullCopyThroughIterator(const sycl::range &shape) { std::vector reference(shape.size()); @@ -373,3 +430,84 @@ TEST_F(AccessorIteratorTest, PartialCopyWithOffset3D) { ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator( sycl::range<3>{9, 8, 7}, sycl::range<3>{3, 4, 5}, sycl::id<3>{3, 2, 1})); } + +TEST_F(AccessorIteratorTest, FullWrite1D) { + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<1>{10}, sycl::range<1>{10}, sycl::id<1>{0})); +} + +TEST_F(AccessorIteratorTest, FullWrite2D) { + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<2>{5, 5}, sycl::range<2>{5, 5}, sycl::id<2>{0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<2>{2, 5}, sycl::range<2>{2, 5}, sycl::id<2>{0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<2>{5, 2}, sycl::range<2>{5, 2}, sycl::id<2>{0, 0})); +} + +TEST_F(AccessorIteratorTest, FullWrite3D) { + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 5, 5}, sycl::id<3>{0, 0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{1, 5, 5}, sycl::range<3>{1, 5, 5}, sycl::id<3>{0, 0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{5, 1, 5}, sycl::range<3>{5, 1, 5}, sycl::id<3>{0, 0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{5, 5, 1}, sycl::range<3>{5, 5, 1}, sycl::id<3>{0, 0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{3, 6, 4}, sycl::range<3>{3, 6, 4}, sycl::id<3>{0, 0, 0})); +} + +TEST_F(AccessorIteratorTest, PartialWriteWithoutOffset1D) { + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<1>{10}, sycl::range<1>{5}, sycl::id<1>{0})); +} + +TEST_F(AccessorIteratorTest, PartialWriteWithoutOffset2D) { + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<2>{5, 5}, sycl::range<2>{3, 3}, sycl::id<2>{0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<2>{2, 5}, sycl::range<2>{1, 3}, sycl::id<2>{0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<2>{5, 2}, sycl::range<2>{3, 1}, sycl::id<2>{0, 0})); +} + +TEST_F(AccessorIteratorTest, PartialWriteWithoutOffset3D) { + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 3}, sycl::id<3>{0, 0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{1, 5, 5}, sycl::range<3>{0, 3, 3}, sycl::id<3>{0, 0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{5, 1, 5}, sycl::range<3>{3, 1, 3}, sycl::id<3>{0, 0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{5, 5, 1}, sycl::range<3>{3, 3, 1}, sycl::id<3>{0, 0, 0})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{3, 6, 4}, sycl::range<3>{1, 3, 2}, sycl::id<3>{0, 0, 0})); +} + +TEST_F(AccessorIteratorTest, PartialWriteWithOffset1D) { + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<1>{10}, sycl::range<1>{5}, sycl::id<1>{3})); +} + +TEST_F(AccessorIteratorTest, PartialWriteWithOffset2D) { + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<2>{5, 5}, sycl::range<2>{3, 3}, sycl::id<2>{1, 1})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<2>{3, 5}, sycl::range<2>{1, 3}, sycl::id<2>{1, 2})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<2>{5, 3}, sycl::range<2>{3, 1}, sycl::id<2>{1, 1})); +} + +TEST_F(AccessorIteratorTest, PartialWriteWithOffset3D) { + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 3}, sycl::id<3>{1, 1, 1})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{3, 5, 5}, sycl::range<3>{0, 3, 3}, sycl::id<3>{1, 2, 2})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{5, 2, 5}, sycl::range<3>{3, 1, 3}, sycl::id<3>{1, 1, 2})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{5, 5, 3}, sycl::range<3>{3, 3, 1}, sycl::id<3>{1, 1, 1})); + ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator( + sycl::range<3>{3, 6, 4}, sycl::range<3>{1, 3, 2}, sycl::id<3>{1, 3, 2})); +} From a6effbb2da67229bf56ce45a9f1a4f8d3ddc6523 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 7 Oct 2022 10:07:45 -0400 Subject: [PATCH 19/21] Revert back from using reserved identifiers --- sycl/include/sycl/accessor.hpp | 12 +- .../include/sycl/detail/accessor_iterator.hpp | 312 +++++++++--------- 2 files changed, 160 insertions(+), 164 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 2b418b5acbd43..a63b6da725646 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -1195,9 +1195,9 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : using reference = DataT &; using const_reference = const DataT &; - using iterator = typename detail::__accessor_iterator; + using iterator = typename detail::accessor_iterator; using const_iterator = - typename detail::__accessor_iterator; + typename detail::accessor_iterator; using difference_type = typename std::iterator_traits::difference_type; @@ -2099,28 +2099,28 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); } iterator begin() const noexcept { - return iterator::__get_begin( + return iterator::getBegin( get_pointer(), detail::convertToArrayOfN(getMemoryRange()), get_range(), get_offset()); } iterator end() const noexcept { - return iterator::__get_end( + return iterator::getEnd( get_pointer(), detail::convertToArrayOfN(getMemoryRange()), get_range(), get_offset()); } const_iterator cbegin() const noexcept { - return const_iterator::__get_begin( + return const_iterator::getBegin( get_pointer(), detail::convertToArrayOfN(getMemoryRange()), get_range(), get_offset()); } const_iterator cend() const noexcept { - return const_iterator::__get_end( + return const_iterator::getEnd( get_pointer(), detail::convertToArrayOfN(getMemoryRange()), get_range(), get_offset()); diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index 39b2f55ac846c..0ecede3bad5f7 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -37,144 +37,142 @@ namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { -template +template class accessor; namespace detail { -template class __accessor_iterator { +template class accessor_iterator { public: using difference_type = std::ptrdiff_t; - using value_type = _DataT; + using value_type = DataT; // FIXME: this should likely include address space - using pointer = _DataT *; - using reference = _DataT &; + using pointer = DataT *; + using reference = DataT &; using iterator_category = std::random_access_iterator_tag; - __accessor_iterator() = default; + accessor_iterator() = default; reference operator*() { - return *(_MDataPtr + __get_absolute_offset_to_buffer()); + return *(MDataPtr + getAbsoluteOffsetToBuffer()); } - __accessor_iterator &operator++() { - ++_MLinearId; + accessor_iterator &operator++() { + ++MLinearId; return *this; } - __accessor_iterator operator++(int) { - auto _Old = *this; + accessor_iterator operator++(int) { + auto Old = *this; ++(*this); - return _Old; + return Old; } - __accessor_iterator &operator--() { - --_MLinearId; + accessor_iterator &operator--() { + --MLinearId; return *this; } - __accessor_iterator operator--(int) { - auto _Old = *this; + accessor_iterator operator--(int) { + auto Old = *this; --(*this); - return _Old; + return Old; } - __accessor_iterator &operator+=(difference_type _N) { - _MLinearId += _N; + accessor_iterator &operator+=(difference_type N) { + MLinearId += N; return *this; } - friend __accessor_iterator operator+(const __accessor_iterator &_Lhs, - difference_type _N) { - auto _Ret = _Lhs; - _Ret += _N; - return _Ret; + friend accessor_iterator operator+(const accessor_iterator &Lhs, + difference_type N) { + auto Ret = Lhs; + Ret += N; + return Ret; } - friend __accessor_iterator operator+(difference_type _N, - const __accessor_iterator &_Rhs) { - auto _Ret = _Rhs; - _Ret += _N; - return _Ret; + friend accessor_iterator operator+(difference_type N, + const accessor_iterator &Rhs) { + auto Ret = Rhs; + Ret += N; + return Ret; } - __accessor_iterator &operator-=(difference_type _N) { - _MLinearId -= _N; + accessor_iterator &operator-=(difference_type N) { + MLinearId -= N; return *this; } - friend __accessor_iterator operator-(__accessor_iterator &_Lhs, - difference_type _N) { - _Lhs -= _N; - return _Lhs; + friend accessor_iterator operator-(accessor_iterator &Lhs, + difference_type N) { + Lhs -= N; + return Lhs; } - reference &operator[](difference_type _N) { - auto _Copy = *this; - _Copy += _N; - return *_Copy; + reference &operator[](difference_type N) { + auto Copy = *this; + Copy += N; + return *Copy; } - bool operator<(const __accessor_iterator &_Other) const { - return _MLinearId < _Other._MLinearId; + bool operator<(const accessor_iterator &Other) const { + return MLinearId < Other.MLinearId; } - bool operator>(const __accessor_iterator &_Other) const { - return _Other < *this; - } + bool operator>(const accessor_iterator &Other) const { return Other < *this; } - bool operator<=(const __accessor_iterator &_Other) const { - return !(*this > _Other); + bool operator<=(const accessor_iterator &Other) const { + return !(*this > Other); } - bool operator>=(const __accessor_iterator &_Other) const { - return !(*this < _Other); + bool operator>=(const accessor_iterator &Other) const { + return !(*this < Other); } - bool operator==(const __accessor_iterator &_Other) const { - return _MLinearId == _Other._MLinearId; + bool operator==(const accessor_iterator &Other) const { + return MLinearId == Other.MLinearId; } - bool operator!=(const __accessor_iterator &_Other) const { - return !(*this == _Other); + bool operator!=(const accessor_iterator &Other) const { + return !(*this == Other); } - difference_type operator-(const __accessor_iterator &_Rhs) { - return _MLinearId - _Rhs._MLinearId; + difference_type operator-(const accessor_iterator &Rhs) { + return MLinearId - Rhs.MLinearId; } private: - template + template friend class sycl::accessor; - _DataT *_MDataPtr = nullptr; + DataT *MDataPtr = nullptr; // Stores a linear id of an accessor's buffer element the iterator points to. // This id is relative to a range accessible through an accessor, i.e. it is // limited by a space with top left corner defiend as accessor::get_offset() // and bottom right corner defined as accesor::get_range(). - size_t _MLinearId = 0; + size_t MLinearId = 0; - // Describes range of linear IDs accessible by the iterator. _MEnd corresponds + // Describes range of linear IDs accessible by the iterator. MEnd corresponds // to ID of en element past the last accessible element of accessors's // buffer. - size_t _MBegin = 0; - size_t _MEnd = 0; + size_t MBegin = 0; + size_t MEnd = 0; // If set to true, then it indicates that accessor has its offset and/or range // set to non-zero, i.e. it is a ranged accessor. - bool _MAccessorIsRanged = false; + bool MAccessorIsRanged = false; // Fields below are used (and changed to be non-zero) only if we deal with // a ranged accessor. // - // TODO: consider making their existance dependable on _Dimensions template + // TODO: consider making their existance dependable on Dimensions template // parameter, because not all of them are needed for all possible dimensions. // Three field below allow us to calculate an absolute offset to an accessor's @@ -196,140 +194,138 @@ template class __accessor_iterator { // . X X X . // . . . . . // - // _MStaticOffset stores a number of elements which precede the first + // MStaticOffset stores a number of elements which precede the first // accessible element, calculated as if the buffer was linearized. - // For the example above, _MStaticOffset would be equal to 6, because + // For the example above, MStaticOffset would be equal to 6, because // there is one full row before the first accessible element and a one more on // the second line. "Static" in the name highlights that this is a constant // element in an equation which calculates an absoulte offset to an accessor's // buffer, it doesn't depend on the current state of the iterator. // - // NOTE: _MStaticOffset is set to 0 in 1D case even if the accessor was + // NOTE: MStaticOffset is set to 0 in 1D case even if the accessor was // created with offset: it is done to further optimize 1D case by - // incorporating that offset into _MLinearId right away. + // incorporating that offset into MLinearId right away. // - // _MPerRowOffset stores a number of _inaccessible_ elements in each - // _accessible_ row. For the example above it would be equal to 2 (leftmost + // MPerRowOffset stores a number of inaccessible_ elements in each + // accessible_ row. For the example above it would be equal to 2 (leftmost // and the rightmost elements of a row). // - // _MPerSliceOffset stores a number of _inaccessible_ elements in each - // _accessible_ slice. Slice here means a single 2D layer in a 3D buffer. For + // MPerSliceOffset stores a number of inaccessible_ elements in each + // accessible_ slice. Slice here means a single 2D layer in a 3D buffer. For // the example above it would be equal to 0, because we are not looking at a // 3D buffer. However, if we had two slices like visualized above, - // _MPerSliceOffset would be equal to 16 (elements on the "perimeter" of the + // MPerSliceOffset would be equal to 16 (elements on the "perimeter" of the // slice, i.e. ones represented as dots (.)). - size_t _MStaticOffset = 0; - size_t _MPerRowOffset = 0; - size_t _MPerSliceOffset = 0; + size_t MStaticOffset = 0; + size_t MPerRowOffset = 0; + size_t MPerSliceOffset = 0; - // Contains a number of _accessible_ elements in a row - size_t _MRowSize = 0; - // Contains a number of _accessible_ elements in a slice - size_t _MSliceSize = 0; + // Contains a number of accessible_ elements in a row + size_t MRowSize = 0; + // Contains a number of accessible_ elements in a slice + size_t MSliceSize = 0; - // _MLinearId stores an offset which is relative to the accessible range of - // the accessor, which means that it could be the case that _MlinearId equal + // MLinearId stores an offset which is relative to the accessible range of + // the accessor, which means that it could be the case that MlinearId equal // to 0 should not correspond to the beginning of the underlying buffer, but // instead should be re-adjusted to account for an offset passed to the // accessor constructor. // // This function performs necessary calculations to make sure that all // access ranges and offsets are taken into account. - size_t __get_absolute_offset_to_buffer() { + size_t getAbsoluteOffsetToBuffer() { // For 1D case, any possible offsets are already incorporated into - // _MLinearId, so 1D is always treated as a non-ranged accessor - if (!_MAccessorIsRanged || _Dimensions == 1) - return _MLinearId; + // MLinearId, so 1D is always treated as a non-ranged accessor + if (!MAccessorIsRanged || Dimensions == 1) + return MLinearId; // Here we need to deal with 2D or 3D ranged accessor. - // _MLinearId points to an element relative to the accessible range. It + // MLinearId points to an element relative to the accessible range. It // should be adjusted to account for elements which are outside of the // accessible range of the accessor. // We start with static offset: that is a number of elements in full rows // and full slices before the first accessible element. - size_t _AbsoluteId = _MLinearId + _MStaticOffset; + size_t AbsoluteId = MLinearId + MStaticOffset; // Then we account for inaccessible elements in each full slice - size_t _Remaining = _MLinearId; - if constexpr (_Dimensions == 3) { - _AbsoluteId += _MPerSliceOffset * (_Remaining / _MSliceSize); - _Remaining %= _MSliceSize; + size_t Remaining = MLinearId; + if constexpr (Dimensions == 3) { + AbsoluteId += MPerSliceOffset * (Remaining / MSliceSize); + Remaining %= MSliceSize; } // Then we account for inaccessible elements in each full row - _AbsoluteId += _MPerRowOffset * (_Remaining / _MRowSize); - _Remaining %= _MRowSize; + AbsoluteId += MPerRowOffset * (Remaining / MRowSize); + Remaining %= MRowSize; - return _AbsoluteId; + return AbsoluteId; } - __accessor_iterator(_DataT *_DataPtr, const range<_Dimensions> &_MemoryRange, - const range<_Dimensions> &_AccessRange, - const id<_Dimensions> &_Offset) - : _MDataPtr(_DataPtr) { - constexpr int _XIndex = _Dimensions - 1; - constexpr int _YIndex = _Dimensions - 2; - (void)_YIndex; - constexpr int _ZIndex = _Dimensions - 3; - (void)_ZIndex; - - if constexpr (_Dimensions > 1) - _MRowSize = _AccessRange[_XIndex]; - if constexpr (_Dimensions > 2) - _MSliceSize = _AccessRange[_YIndex] * _MRowSize; - - if (id<_Dimensions>{} != _Offset) - _MAccessorIsRanged = true; + accessor_iterator(DataT *DataPtr, const range &MemoryRange, + const range &AccessRange, + const id &Offset) + : MDataPtr(DataPtr) { + constexpr int XIndex = Dimensions - 1; + constexpr int YIndex = Dimensions - 2; + (void)YIndex; + constexpr int ZIndex = Dimensions - 3; + (void)ZIndex; + + if constexpr (Dimensions > 1) + MRowSize = AccessRange[XIndex]; + if constexpr (Dimensions > 2) + MSliceSize = AccessRange[YIndex] * MRowSize; + + if (id{} != Offset) + MAccessorIsRanged = true; else { - for (size_t _I = 0; _I < _Dimensions; ++_I) - if (_AccessRange[_I] != _MemoryRange[_I]) - _MAccessorIsRanged = true; + for (size_t I = 0; I < Dimensions; ++I) + if (AccessRange[I] != MemoryRange[I]) + MAccessorIsRanged = true; } - if (_MAccessorIsRanged) { - if constexpr (_Dimensions > 2) { - _MStaticOffset += - _MemoryRange[_XIndex] * _MemoryRange[_YIndex] * _Offset[_ZIndex]; - _MPerSliceOffset = - _MemoryRange[_XIndex] * _MemoryRange[_YIndex] - _MSliceSize; + if (MAccessorIsRanged) { + if constexpr (Dimensions > 2) { + MStaticOffset += + MemoryRange[XIndex] * MemoryRange[YIndex] * Offset[ZIndex]; + MPerSliceOffset = + MemoryRange[XIndex] * MemoryRange[YIndex] - MSliceSize; } - if constexpr (_Dimensions > 1) { + if constexpr (Dimensions > 1) { // Elements in fully inaccessible rows - _MStaticOffset += _MemoryRange[_XIndex] * _Offset[_YIndex]; - _MPerRowOffset = _MemoryRange[_XIndex] - _MRowSize; + MStaticOffset += MemoryRange[XIndex] * Offset[YIndex]; + MPerRowOffset = MemoryRange[XIndex] - MRowSize; } // Elements from the first accessible row - if constexpr (_Dimensions == 1) - // To further optimize 1D case, offset is already included into _Begin - _MBegin = _Offset[_XIndex]; + if constexpr (Dimensions == 1) + // To further optimize 1D case, offset is already included into Begin + MBegin = Offset[XIndex]; else - _MStaticOffset += _Offset[_XIndex]; + MStaticOffset += Offset[XIndex]; } - _MEnd = _MBegin + _AccessRange.size(); + MEnd = MBegin + AccessRange.size(); } - static __accessor_iterator __get_begin(_DataT *_DataPtr, - const range<_Dimensions> &_MemoryRange, - const range<_Dimensions> &_AccessRange, - const id<_Dimensions> &_Offset) { - auto _It = - __accessor_iterator(_DataPtr, _MemoryRange, _AccessRange, _Offset); - _It._MLinearId = _It._MBegin; - return _It; + static accessor_iterator getBegin(DataT *DataPtr, + const range &MemoryRange, + const range &AccessRange, + const id &Offset) { + auto It = accessor_iterator(DataPtr, MemoryRange, AccessRange, Offset); + It.MLinearId = It.MBegin; + return It; } - static __accessor_iterator __get_end(_DataT *_DataPtr, - const range<_Dimensions> &_MemoryRange, - const range<_Dimensions> &_AccessRange, - const id<_Dimensions> &_Offset) { - auto _It = - __accessor_iterator(_DataPtr, _MemoryRange, _AccessRange, _Offset); - _It._MLinearId = _It._MEnd; - return _It; + static accessor_iterator getEnd(DataT *DataPtr, + const range &MemoryRange, + const range &AccessRange, + const id &Offset) { + auto It = accessor_iterator(DataPtr, MemoryRange, AccessRange, Offset); + It.MLinearId = It.MEnd; + return It; } public: @@ -337,21 +333,21 @@ template class __accessor_iterator { // Could be useful for debugging, but not a part of the official API, // therefore only available in builds with assertions enabled. friend std::ostream &operator<<(std::ostream &os, - const __accessor_iterator &it) { - os << "__accessor_iterator {\n"; - os << "\t_MLinearId: " << it._MLinearId << "\n"; - os << "\t_MEnd: " << it._MEnd << "\n"; - os << "\t_MStaticOffset: " << it._MStaticOffset << "\n"; - os << "\t_MPerRowOffset: " << it._MPerRowOffset << "\n"; - os << "\t_MPerSliceOffset: " << it._MPerSliceOffset << "\n"; - os << "\t_MRowSize: " << it._MRowSize << "\n"; - os << "\t_MSliceSize: " << it._MSliceSize << "\n"; - os << "\t_MAccessorIsRanged: " << it._MAccessorIsRanged << "\n"; + const accessor_iterator &it) { + os << "accessor_iterator {\n"; + os << "\tMLinearId: " << it.MLinearId << "\n"; + os << "\tMEnd: " << it.MEnd << "\n"; + os << "\tMStaticOffset: " << it.MStaticOffset << "\n"; + os << "\tMPerRowOffset: " << it.MPerRowOffset << "\n"; + os << "\tMPerSliceOffset: " << it.MPerSliceOffset << "\n"; + os << "\tMRowSize: " << it.MRowSize << "\n"; + os << "\tMSliceSize: " << it.MSliceSize << "\n"; + os << "\tMAccessorIsRanged: " << it.MAccessorIsRanged << "\n"; os << "}"; return os; } #endif // NDEBUG }; } // namespace detail -} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // __SYCL_INLINE_VER_NAMESPACE(V1) } // namespace sycl From 600916bc7cb28ee1245c18cbbf873ab66bd04cfa Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 7 Oct 2022 10:09:27 -0400 Subject: [PATCH 20/21] Fix clang-format --- sycl/include/sycl/detail/accessor_iterator.hpp | 18 ++++++++---------- 1 file changed, 8 insertions(+), 10 deletions(-) diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index 0ecede3bad5f7..d91e8c568fcf0 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -55,9 +55,7 @@ template class accessor_iterator { accessor_iterator() = default; - reference operator*() { - return *(MDataPtr + getAbsoluteOffsetToBuffer()); - } + reference operator*() { return *(MDataPtr + getAbsoluteOffsetToBuffer()); } accessor_iterator &operator++() { ++MLinearId; @@ -311,18 +309,18 @@ template class accessor_iterator { } static accessor_iterator getBegin(DataT *DataPtr, - const range &MemoryRange, - const range &AccessRange, - const id &Offset) { + const range &MemoryRange, + const range &AccessRange, + const id &Offset) { auto It = accessor_iterator(DataPtr, MemoryRange, AccessRange, Offset); It.MLinearId = It.MBegin; return It; } static accessor_iterator getEnd(DataT *DataPtr, - const range &MemoryRange, - const range &AccessRange, - const id &Offset) { + const range &MemoryRange, + const range &AccessRange, + const id &Offset) { auto It = accessor_iterator(DataPtr, MemoryRange, AccessRange, Offset); It.MLinearId = It.MEnd; return It; @@ -349,5 +347,5 @@ template class accessor_iterator { #endif // NDEBUG }; } // namespace detail -} // __SYCL_INLINE_VER_NAMESPACE(V1) +} // __SYCL_INLINE_VER_NAMESPACE(_V1) } // namespace sycl From 78215dcdbcb2905c3ee9ca4b0bc0e2bebf82ffcb Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Fri, 7 Oct 2022 10:16:09 -0400 Subject: [PATCH 21/21] Return some undescores from comments back --- sycl/include/sycl/detail/accessor_iterator.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp index d91e8c568fcf0..244350b5cb7fb 100644 --- a/sycl/include/sycl/detail/accessor_iterator.hpp +++ b/sycl/include/sycl/detail/accessor_iterator.hpp @@ -204,12 +204,12 @@ template class accessor_iterator { // created with offset: it is done to further optimize 1D case by // incorporating that offset into MLinearId right away. // - // MPerRowOffset stores a number of inaccessible_ elements in each - // accessible_ row. For the example above it would be equal to 2 (leftmost + // MPerRowOffset stores a number of _inaccessible_ elements in each + // _accessible_ row. For the example above it would be equal to 2 (leftmost // and the rightmost elements of a row). // - // MPerSliceOffset stores a number of inaccessible_ elements in each - // accessible_ slice. Slice here means a single 2D layer in a 3D buffer. For + // MPerSliceOffset stores a number of _inaccessible_ elements in each + // _accessible_ slice. Slice here means a single 2D layer in a 3D buffer. For // the example above it would be equal to 0, because we are not looking at a // 3D buffer. However, if we had two slices like visualized above, // MPerSliceOffset would be equal to 16 (elements on the "perimeter" of the @@ -219,9 +219,9 @@ template class accessor_iterator { size_t MPerRowOffset = 0; size_t MPerSliceOffset = 0; - // Contains a number of accessible_ elements in a row + // Contains a number of _accessible_ elements in a row size_t MRowSize = 0; - // Contains a number of accessible_ elements in a slice + // Contains a number of _accessible_ elements in a slice size_t MSliceSize = 0; // MLinearId stores an offset which is relative to the accessible range of