From f905d91e941479e841c98fe33ac2ca61574e7278 Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Thu, 27 Jan 2022 13:05:10 +0000 Subject: [PATCH 1/9] Rename Iterator to Memory Access Strategy - rename was necessary to avoid confusions with fancy iterator - rename and flat namespace of the mem access strategy - move related files to new folder --- include/vikunja/access/BaseStrategy.hpp | 127 +++++++ .../access/PolicyBasedBlockStrategy.hpp | 317 +++++++++++++++++ include/vikunja/mem/iterator/BaseIterator.hpp | 135 -------- .../mem/iterator/PolicyBasedBlockIterator.hpp | 326 ------------------ .../reduce/detail/BlockThreadReduceKernel.hpp | 6 +- include/vikunja/reduce/reduce.hpp | 16 +- .../detail/BlockThreadTransformKernel.hpp | 24 +- include/vikunja/transform/transform.hpp | 16 +- test/include/vikunja/test/utility.hpp | 6 +- test/unit/iterator/src/Iterator.cpp | 22 +- 10 files changed, 492 insertions(+), 503 deletions(-) create mode 100644 include/vikunja/access/BaseStrategy.hpp create mode 100644 include/vikunja/access/PolicyBasedBlockStrategy.hpp delete mode 100644 include/vikunja/mem/iterator/BaseIterator.hpp delete mode 100644 include/vikunja/mem/iterator/PolicyBasedBlockIterator.hpp diff --git a/include/vikunja/access/BaseStrategy.hpp b/include/vikunja/access/BaseStrategy.hpp new file mode 100644 index 0000000..9efc18c --- /dev/null +++ b/include/vikunja/access/BaseStrategy.hpp @@ -0,0 +1,127 @@ +/* Copyright 2022 Hauke Mewes, Jonas Schenke, Simeon Ehrig + * + * This file is part of vikunja. + * + * This Source Code Form is subject to the terms of the Mozilla Public + * License, v. 2.0. If a copy of the MPL was not distributed with this + * file, You can obtain one at http://mozilla.org/MPL/2.0/. + */ + +#pragma once + +#include + +namespace vikunja::MemAccess +{ + //! An iterator base class. + //! + //! \tparam T The type. + //! \tparam TIterator The iterator type (standard is T*) + template + class BaseStrategy + { + protected: + TIterator const mData; // The underlying iterator must be const as we pass it by ref + uint64_t mIndex; + const uint64_t mMaximum; + + public: + //----------------------------------------------------------------------------- + //! Constructor. + //! + //! \param data A pointer to the data. + //! \param index The index. + //! \param maximum The first index outside of the iterator memory. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE BaseStrategy(TIterator const& data, uint64_t index, uint64_t maximum) + : mData(data) + , mIndex(index) + , mMaximum(maximum) + { + } + + //----------------------------------------------------------------------------- + //! Constructor. + //! + //! \param other The other iterator object. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE BaseStrategy(const BaseStrategy& other) = default; + + //----------------------------------------------------------------------------- + //! Compare operator. + //! + //! \param other The other object. + //! + //! Returns true if objects are equal and false otherwise. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator==(const BaseStrategy& other) const -> bool + { + return (this->mData == other.mData) && (this->mIndex == other.mIndex) + && (this->mMaximum == other.mMaximum); + } + + //----------------------------------------------------------------------------- + //! Compare operator. + //! + //! \param other The other object. + //! + //! Returns false if objects are equal and true otherwise. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator!=(const BaseStrategy& other) const -> bool + { + return !operator==(other); + } + + //----------------------------------------------------------------------------- + //! Compare operator. + //! + //! \param other The other object. + //! + //! Returns false if the other object is equal or smaller and true + //! otherwise. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator<(const BaseStrategy& other) const -> bool + { + return mIndex < other.mIndex; + } + + //----------------------------------------------------------------------------- + //! Compare operator. + //! + //! \param other The other object. + //! + //! Returns false if the other object is equal or bigger and true otherwise. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator>(const BaseStrategy& other) const -> bool + { + return mIndex > other.mIndex; + } + + //----------------------------------------------------------------------------- + //! Compare operator. + //! + //! \param other The other object. + //! + //! Returns true if the other object is equal or bigger and false otherwise. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator<=(const BaseStrategy& other) const -> bool + { + return mIndex <= other.mIndex; + } + + //----------------------------------------------------------------------------- + //! Compare operator. + //! + //! \param other The other object. + //! + //! Returns true if the other object is equal or smaller and false + //! otherwise. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator>=(const BaseStrategy& other) const -> bool + { + return mIndex >= other.mIndex; + } + + //----------------------------------------------------------------------------- + //! Returns the current element. + //! + //! Returns a reference to the current index. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator*() -> decltype(*(mData + mIndex))& + { + return *(mData + mIndex); + } + }; + +} // namespace vikunja::MemAccess diff --git a/include/vikunja/access/PolicyBasedBlockStrategy.hpp b/include/vikunja/access/PolicyBasedBlockStrategy.hpp new file mode 100644 index 0000000..ab4418a --- /dev/null +++ b/include/vikunja/access/PolicyBasedBlockStrategy.hpp @@ -0,0 +1,317 @@ +/* Copyright 2022 Hauke Mewes, Simeon Ehrig + * + * This file is part of vikunja. + * + * This Source Code Form is subject to the terms of the Mozilla Public + * License, v. 2.0. If a copy of the MPL was not distributed with this + * file, You can obtain one at http://mozilla.org/MPL/2.0/. + */ + +#pragma once + +#include "BaseStrategy.hpp" + +#include + +namespace vikunja::MemAccess +{ + /** + * A policy based iterator that splits the data into chunks. Depending on the memory access policy, + * these chunks can access the data sequential or in a striding pattern. + * @tparam MemAccessPolicy The memory access policy to use. + * @tparam TAcc The alpaka accelerator type. + * @tparam TIterator The underlying iterator. + * + * The memory access policy should provide three values: + * - The startIndex of the iterator, which is the first index to use. + * - The endIndex of the iterator, which is the last index to use. + * - The stepSize of the iterator, which tells how far the iterator should move. + */ + template + class PolicyBasedBlockStrategy : public BaseStrategy + { + private: + uint64_t mStep; /**< The step size of this iterator. */ + public: + /** + * Create a policy based block iterator + * @param data The data to iterate over. + * @param acc The accelerator type to use. + * @param problemSize The size of the original iterator. + * @param blockSize The size of the blocks. + */ + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE + PolicyBasedBlockStrategy(TIterator const& data, TAcc const& acc, uint64_t problemSize, uint64_t blockSize) + : BaseStrategy( + data, + MemAccessPolicy::getStartIndex(acc, problemSize, blockSize), + MemAccessPolicy::getEndIndex(acc, problemSize, blockSize)) + , mStep(MemAccessPolicy::getStepSize(acc, problemSize, blockSize)) + { + } + + /** + * Default copy constructor. + * @param other To be copied. + */ + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE PolicyBasedBlockStrategy(const PolicyBasedBlockStrategy& other) = default; + + //----------------------------------------------------------------------------- + //! Returns the iterator for the last item. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto end() const -> PolicyBasedBlockStrategy + { + PolicyBasedBlockStrategy ret(*this); + ret.mIndex = this->mMaximum; + return ret; + } + + //----------------------------------------------------------------------------- + //! Increments the internal pointer to the next one and returns this + //! element. + //! + //! Returns a reference to the next index. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator++() -> PolicyBasedBlockStrategy& + { + this->mIndex += this->mStep; + return *this; + } + + //----------------------------------------------------------------------------- + //! Returns the current element and increments the internal pointer to the + //! next one. + //! + //! Returns a reference to the current index. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator++(int) -> PolicyBasedBlockStrategy + { + auto ret(*this); + this->mIndex += this->mStep; + return ret; + } + + //----------------------------------------------------------------------------- + //! Decrements the internal pointer to the previous one and returns the this + //! element. + //! + //! Returns a reference to the previous index. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator--() -> PolicyBasedBlockStrategy& + { + this->mIndex -= this->mStep; + return *this; + } + + //----------------------------------------------------------------------------- + //! Returns the current element and decrements the internal pointer to the + //! previous one. + //! + //! Returns a reference to the current index. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator--(int) -> PolicyBasedBlockStrategy + { + auto ret(*this); + this->mIndex -= this->mStep; + return ret; + } + + //----------------------------------------------------------------------------- + //! Returns the index + a supplied offset. + //! + //! \param n The offset. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator+(uint64_t n) const -> PolicyBasedBlockStrategy + { + auto ret(*this); + ret.mIndex += n * mStep; + return ret; + } + + //----------------------------------------------------------------------------- + //! Returns the index - a supplied offset. + //! + //! \param n The offset. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator-(uint64_t n) const -> PolicyBasedBlockStrategy + { + auto ret(*this); + ret.mIndex -= n * mStep; + return ret; + } + + //----------------------------------------------------------------------------- + //! Addition assignment. + //! + //! \param offset The offset. + //! + //! Returns the current object offset by the offset. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator+=(uint64_t offset) -> PolicyBasedBlockStrategy& + { + this->mIndex += offset * this->mStep; + return *this; + } + + //----------------------------------------------------------------------------- + //! Substraction assignment. + //! + //! \param offset The offset. + //! + //! Returns the current object offset by the offset. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator-=(uint64_t offset) -> PolicyBasedBlockStrategy& + { + this->mIndex -= offset * this->mStep; + return *this; + } + }; + + namespace policies + { + /** + * A memory policy for the PolicyBlockBasedIterator that provides grid striding memory access. + */ + struct GridStridingMemAccessPolicy + { + template + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto getStartIndex( + TAcc const& acc, + TIdx const& problemSize __attribute__((unused)), + TIdx const& blockSize __attribute__((unused))) -> TIdx const + { + constexpr TIdx xIndex = alpaka::Dim::value - 1u; + return alpaka::getIdx(acc)[xIndex]; + } + + template + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto getEndIndex( + TAcc const& acc, + TIdx const& problemSize, + TIdx const& blockSize __attribute__((unused))) -> TIdx const + { + return problemSize; + } + + template + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto getStepSize( + TAcc const& acc, + TIdx const& problemSize __attribute__((unused)), + TIdx const& blockSize) -> TIdx const + { + constexpr TIdx xIndex = alpaka::Dim::value - 1u; + auto gridDimension = alpaka::getWorkDiv(acc)[xIndex]; + return gridDimension * blockSize; + } + + template + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto isValidThreadResult( + TAcc const& acc, + TIdx const& problemSize, + TIdx const& blockSize __attribute__((unused))) -> bool const + { + constexpr TIdx xIndex = alpaka::Dim::value - 1u; + auto threadIndex = (alpaka::getIdx(acc)[xIndex]); + return threadIndex < problemSize; + } + + static constexpr bool isThreadOrderCompliant = true; + + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr auto getName() -> char* + { + return const_cast("GridStridingMemAccessPolicy"); + } + }; + + /** + * A memory access policy for the PolicyBlockBasedIterator that provides linear memory access. + */ + struct LinearMemAccessPolicy + { + template + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto getStartIndex( + TAcc const& acc, + TIdx const& problemSize, + TIdx const& blockSize) -> TIdx const + { + constexpr TIdx xIndex = alpaka::Dim::value - 1u; + auto gridDimension = alpaka::getWorkDiv(acc)[xIndex]; + auto indexInBlock = alpaka::getIdx(acc)[xIndex]; + auto gridSize = gridDimension * blockSize; + // TODO: catch overflow + return (problemSize * indexInBlock) / gridSize; + } + + template + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto getEndIndex( + TAcc const& acc, + TIdx const& problemSize, + TIdx const& blockSize) -> TIdx const + { + constexpr TIdx xIndex = alpaka::Dim::value - 1u; + auto gridDimension = alpaka::getWorkDiv(acc)[xIndex]; + auto indexInBlock = alpaka::getIdx(acc)[xIndex]; + auto gridSize = gridDimension * blockSize; + // TODO: catch overflow + return (problemSize * indexInBlock + problemSize) / gridSize; + } + + template + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr auto getStepSize( + TAcc const& acc __attribute__((unused)), + TIdx const& problemSize __attribute__((unused)), + TIdx const& blockSize __attribute__((unused))) -> TIdx const + { + return 1; + } + + template + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr auto isValidThreadResult( + TAcc const& acc __attribute__((unused)), + TIdx const& problemSize __attribute__((unused)), + TIdx const& blockSize __attribute__((unused))) -> bool const + { + return true; + } + + static constexpr bool isThreadOrderCompliant = false; + + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr auto getName() -> char* + { + return const_cast("LinearMemAccessPolicy"); + } + }; + } // namespace policies + + namespace traits + { + /** + * The memory access policy getter trait by platform. + * @tparam TAcc The platform type. + * @tparam TSfinae + */ + template + struct GetMemAccessPolicyByPltf + { + }; + + /** + * On cpu, default memory access is linear. + */ + template<> + struct GetMemAccessPolicyByPltf + { + using type = policies::LinearMemAccessPolicy; + }; + +#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) || defined(ALPAKA_ACC_GPU_HIP_ENABLED) + /** + * On cuda, default memory access is grid striding. + */ + template<> + struct GetMemAccessPolicyByPltf + { + using type = policies::GridStridingMemAccessPolicy; + }; +#endif // (ALPAKA_ACC_GPU_CUDA_ENABLED) || defined(ALPAKA_ACC_GPU_HIP_ENABLED) + + } // namespace traits + + /** + * Shortcut to derive memory access policy from accelerator. + */ + template + using MemAccessPolicy = typename traits::GetMemAccessPolicyByPltf>>::type; + +} // namespace vikunja::MemAccess diff --git a/include/vikunja/mem/iterator/BaseIterator.hpp b/include/vikunja/mem/iterator/BaseIterator.hpp deleted file mode 100644 index 1a4a1d3..0000000 --- a/include/vikunja/mem/iterator/BaseIterator.hpp +++ /dev/null @@ -1,135 +0,0 @@ -/* Copyright 2021 Hauke Mewes, Jonas Schenke - * - * This file is part of vikunja. - * - * This Source Code Form is subject to the terms of the Mozilla Public - * License, v. 2.0. If a copy of the MPL was not distributed with this - * file, You can obtain one at http://mozilla.org/MPL/2.0/. - */ - -#pragma once - -#include - -namespace vikunja -{ - namespace mem - { - namespace iterator - { - // TODO: this class is from Jonas Schenke - - //! An iterator base class. - //! - //! \tparam T The type. - //! \tparam TIterator The iterator type (standard is T*) - template - class BaseIterator - { - protected: - TIterator const mData; // The underlying iterator must be const as we pass it by ref - uint64_t mIndex; - const uint64_t mMaximum; - - public: - //----------------------------------------------------------------------------- - //! Constructor. - //! - //! \param data A pointer to the data. - //! \param index The index. - //! \param maximum The first index outside of the iterator memory. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE - BaseIterator(TIterator const& data, uint64_t index, uint64_t maximum) - : mData(data) - , mIndex(index) - , mMaximum(maximum) - { - } - - //----------------------------------------------------------------------------- - //! Constructor. - //! - //! \param other The other iterator object. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE BaseIterator(const BaseIterator& other) = default; - - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns true if objects are equal and false otherwise. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator==(const BaseIterator& other) const -> bool - { - return (this->mData == other.mData) && (this->mIndex == other.mIndex) - && (this->mMaximum == other.mMaximum); - } - - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns false if objects are equal and true otherwise. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator!=(const BaseIterator& other) const -> bool - { - return !operator==(other); - } - - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns false if the other object is equal or smaller and true - //! otherwise. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator<(const BaseIterator& other) const -> bool - { - return mIndex < other.mIndex; - } - - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns false if the other object is equal or bigger and true otherwise. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator>(const BaseIterator& other) const -> bool - { - return mIndex > other.mIndex; - } - - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns true if the other object is equal or bigger and false otherwise. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator<=(const BaseIterator& other) const -> bool - { - return mIndex <= other.mIndex; - } - - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns true if the other object is equal or smaller and false - //! otherwise. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator>=(const BaseIterator& other) const -> bool - { - return mIndex >= other.mIndex; - } - - //----------------------------------------------------------------------------- - //! Returns the current element. - //! - //! Returns a reference to the current index. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator*() -> decltype(*(mData + mIndex))& - { - return *(mData + mIndex); - } - }; - } // namespace iterator - } // namespace mem -} // namespace vikunja diff --git a/include/vikunja/mem/iterator/PolicyBasedBlockIterator.hpp b/include/vikunja/mem/iterator/PolicyBasedBlockIterator.hpp deleted file mode 100644 index 3e5f70b..0000000 --- a/include/vikunja/mem/iterator/PolicyBasedBlockIterator.hpp +++ /dev/null @@ -1,326 +0,0 @@ -/* Copyright 2021 Hauke Mewes - * - * This file is part of vikunja. - * - * This Source Code Form is subject to the terms of the Mozilla Public - * License, v. 2.0. If a copy of the MPL was not distributed with this - * file, You can obtain one at http://mozilla.org/MPL/2.0/. - */ - -#pragma once - -#include "BaseIterator.hpp" - -#include - -namespace vikunja -{ - namespace mem - { - namespace iterator - { - /** - * A policy based iterator that splits the data into chunks. Depending on the memory access policy, - * these chunks can access the data sequential or in a striding pattern. - * @tparam MemAccessPolicy The memory access policy to use. - * @tparam TAcc The alpaka accelerator type. - * @tparam TIterator The underlying iterator. - * - * The memory access policy should provide three values: - * - The startIndex of the iterator, which is the first index to use. - * - The endIndex of the iterator, which is the last index to use. - * - The stepSize of the iterator, which tells how far the iterator should move. - */ - template - class PolicyBasedBlockIterator : public BaseIterator - { - private: - uint64_t mStep; /**< The step size of this iterator. */ - public: - /** - * Create a policy based block iterator - * @param data The data to iterate over. - * @param acc The accelerator type to use. - * @param problemSize The size of the original iterator. - * @param blockSize The size of the blocks. - */ - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE PolicyBasedBlockIterator( - TIterator const& data, - TAcc const& acc, - uint64_t problemSize, - uint64_t blockSize) - : BaseIterator( - data, - MemAccessPolicy::getStartIndex(acc, problemSize, blockSize), - MemAccessPolicy::getEndIndex(acc, problemSize, blockSize)) - , mStep(MemAccessPolicy::getStepSize(acc, problemSize, blockSize)) - { - } - - /** - * Default copy constructor. - * @param other To be copied. - */ - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE PolicyBasedBlockIterator(const PolicyBasedBlockIterator& other) - = default; - - //----------------------------------------------------------------------------- - //! Returns the iterator for the last item. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto end() const -> PolicyBasedBlockIterator - { - PolicyBasedBlockIterator ret(*this); - ret.mIndex = this->mMaximum; - return ret; - } - - //----------------------------------------------------------------------------- - //! Increments the internal pointer to the next one and returns this - //! element. - //! - //! Returns a reference to the next index. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator++() -> PolicyBasedBlockIterator& - { - this->mIndex += this->mStep; - return *this; - } - - //----------------------------------------------------------------------------- - //! Returns the current element and increments the internal pointer to the - //! next one. - //! - //! Returns a reference to the current index. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator++(int) -> PolicyBasedBlockIterator - { - auto ret(*this); - this->mIndex += this->mStep; - return ret; - } - - //----------------------------------------------------------------------------- - //! Decrements the internal pointer to the previous one and returns the this - //! element. - //! - //! Returns a reference to the previous index. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator--() -> PolicyBasedBlockIterator& - { - this->mIndex -= this->mStep; - return *this; - } - - //----------------------------------------------------------------------------- - //! Returns the current element and decrements the internal pointer to the - //! previous one. - //! - //! Returns a reference to the current index. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator--(int) -> PolicyBasedBlockIterator - { - auto ret(*this); - this->mIndex -= this->mStep; - return ret; - } - - //----------------------------------------------------------------------------- - //! Returns the index + a supplied offset. - //! - //! \param n The offset. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator+(uint64_t n) const -> PolicyBasedBlockIterator - { - auto ret(*this); - ret.mIndex += n * mStep; - return ret; - } - - //----------------------------------------------------------------------------- - //! Returns the index - a supplied offset. - //! - //! \param n The offset. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator-(uint64_t n) const -> PolicyBasedBlockIterator - { - auto ret(*this); - ret.mIndex -= n * mStep; - return ret; - } - - //----------------------------------------------------------------------------- - //! Addition assignment. - //! - //! \param offset The offset. - //! - //! Returns the current object offset by the offset. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator+=(uint64_t offset) -> PolicyBasedBlockIterator& - { - this->mIndex += offset * this->mStep; - return *this; - } - - //----------------------------------------------------------------------------- - //! Substraction assignment. - //! - //! \param offset The offset. - //! - //! Returns the current object offset by the offset. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator-=(uint64_t offset) -> PolicyBasedBlockIterator& - { - this->mIndex -= offset * this->mStep; - return *this; - } - }; - - namespace policies - { - /** - * A memory policy for the PolicyBlockBasedIterator that provides grid striding memory access. - */ - struct GridStridingMemAccessPolicy - { - template - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto getStartIndex( - TAcc const& acc, - TIdx const& problemSize __attribute__((unused)), - TIdx const& blockSize __attribute__((unused))) -> TIdx const - { - constexpr TIdx xIndex = alpaka::Dim::value - 1u; - return alpaka::getIdx(acc)[xIndex]; - } - - template - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto getEndIndex( - TAcc const& acc, - TIdx const& problemSize, - TIdx const& blockSize __attribute__((unused))) -> TIdx const - { - return problemSize; - } - - template - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto getStepSize( - TAcc const& acc, - TIdx const& problemSize __attribute__((unused)), - TIdx const& blockSize) -> TIdx const - { - constexpr TIdx xIndex = alpaka::Dim::value - 1u; - auto gridDimension = alpaka::getWorkDiv(acc)[xIndex]; - return gridDimension * blockSize; - } - - template - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto isValidThreadResult( - TAcc const& acc, - TIdx const& problemSize, - TIdx const& blockSize __attribute__((unused))) -> bool const - { - constexpr TIdx xIndex = alpaka::Dim::value - 1u; - auto threadIndex = (alpaka::getIdx(acc)[xIndex]); - return threadIndex < problemSize; - } - - static constexpr bool isThreadOrderCompliant = true; - - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr auto getName() -> char* - { - return const_cast("GridStridingMemAccessPolicy"); - } - }; - - /** - * A memory access policy for the PolicyBlockBasedIterator that provides linear memory access. - */ - struct LinearMemAccessPolicy - { - template - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto getStartIndex( - TAcc const& acc, - TIdx const& problemSize, - TIdx const& blockSize) -> TIdx const - { - constexpr TIdx xIndex = alpaka::Dim::value - 1u; - auto gridDimension = alpaka::getWorkDiv(acc)[xIndex]; - auto indexInBlock = alpaka::getIdx(acc)[xIndex]; - auto gridSize = gridDimension * blockSize; - // TODO: catch overflow - return (problemSize * indexInBlock) / gridSize; - } - - template - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static auto getEndIndex( - TAcc const& acc, - TIdx const& problemSize, - TIdx const& blockSize) -> TIdx const - { - constexpr TIdx xIndex = alpaka::Dim::value - 1u; - auto gridDimension = alpaka::getWorkDiv(acc)[xIndex]; - auto indexInBlock = alpaka::getIdx(acc)[xIndex]; - auto gridSize = gridDimension * blockSize; - // TODO: catch overflow - return (problemSize * indexInBlock + problemSize) / gridSize; - } - - template - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr auto getStepSize( - TAcc const& acc __attribute__((unused)), - TIdx const& problemSize __attribute__((unused)), - TIdx const& blockSize __attribute__((unused))) -> TIdx const - { - return 1; - } - - template - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr auto isValidThreadResult( - TAcc const& acc __attribute__((unused)), - TIdx const& problemSize __attribute__((unused)), - TIdx const& blockSize __attribute__((unused))) -> bool const - { - return true; - } - - static constexpr bool isThreadOrderCompliant = false; - - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE static constexpr auto getName() -> char* - { - return const_cast("LinearMemAccessPolicy"); - } - }; - } // namespace policies - - namespace traits - { - /** - * The memory access policy getter trait by platform. - * @tparam TAcc The platform type. - * @tparam TSfinae - */ - template - struct GetMemAccessPolicyByPltf - { - }; - - /** - * On cpu, default memory access is linear. - */ - template<> - struct GetMemAccessPolicyByPltf - { - using type = policies::LinearMemAccessPolicy; - }; - -#if defined(ALPAKA_ACC_GPU_CUDA_ENABLED) || defined(ALPAKA_ACC_GPU_HIP_ENABLED) - /** - * On cuda, default memory access is grid striding. - */ - template<> - struct GetMemAccessPolicyByPltf - { - using type = policies::GridStridingMemAccessPolicy; - }; -#endif // (ALPAKA_ACC_GPU_CUDA_ENABLED) || defined(ALPAKA_ACC_GPU_HIP_ENABLED) - - } // namespace traits - - /** - * Shortcut to derive memory access policy from accelerator. - */ - template - using MemAccessPolicy = typename traits::GetMemAccessPolicyByPltf>>::type; - } // namespace iterator - } // namespace mem -} // namespace vikunja diff --git a/include/vikunja/reduce/detail/BlockThreadReduceKernel.hpp b/include/vikunja/reduce/detail/BlockThreadReduceKernel.hpp index cafbb89..dde2c77 100644 --- a/include/vikunja/reduce/detail/BlockThreadReduceKernel.hpp +++ b/include/vikunja/reduce/detail/BlockThreadReduceKernel.hpp @@ -1,4 +1,4 @@ -/* Copyright 2021 Hauke Mewes +/* Copyright 2022 Hauke Mewes, Simeon Ehrig * * This file is part of vikunja. * @@ -9,7 +9,7 @@ #pragma once -#include +#include #include @@ -104,7 +104,7 @@ namespace vikunja using MemPolicy = TMemAccessPolicy; // Create an iterator with the specified memory access policy that wraps the input iterator. - vikunja::mem::iterator::PolicyBasedBlockIterator iter( + vikunja::MemAccess::PolicyBasedBlockStrategy iter( source, acc, n, diff --git a/include/vikunja/reduce/reduce.hpp b/include/vikunja/reduce/reduce.hpp index 906ed43..6537116 100644 --- a/include/vikunja/reduce/reduce.hpp +++ b/include/vikunja/reduce/reduce.hpp @@ -1,4 +1,4 @@ -/* Copyright 2021 Hauke Mewes, Simeon Ehrig +/* Copyright 2022 Hauke Mewes, Simeon Ehrig * * This file is part of vikunja. * @@ -9,7 +9,7 @@ #pragma once -#include +#include #include #include #include @@ -54,7 +54,7 @@ namespace vikunja * @tparam WorkDivPolicy The working division policy. Defaults to a templated value depending on the * accelerator. For the API of this, see workdiv/BlockBasedWorkDiv.hpp * @tparam MemAccessPolicy The memory access policy. Defaults to a templated value depending on the - * accelerator. For the API of this, see mem/iterator/PolicyBasedBlockIterator + * accelerator. For the API of this, see vikunja::MemAccess::PolicyBasedBlockStrategy * @tparam TTransformFunc Type of the transform operator. * @tparam TReduceFunc Type of the reduce operator. * @tparam TInputIterator Type of the input iterator. Should be a pointer-like type. @@ -77,7 +77,7 @@ namespace vikunja template< typename TAcc, typename WorkDivPolicy = vikunja::workdiv::BlockBasedPolicy, - typename MemAccessPolicy = vikunja::mem::iterator::MemAccessPolicy, + typename MemAccessPolicy = vikunja::MemAccess::MemAccessPolicy, typename TTransformFunc, typename TReduceFunc, typename TInputIterator, @@ -219,7 +219,7 @@ namespace vikunja * @tparam WorkDivPolicy The working division policy. Defaults to a templated value depending on the * accelerator. For the API of this, see workdiv/BlockBasedWorkDiv.hpp * @tparam MemAccessPolicy The memory access policy. Defaults to a templated value depending on the - * accelerator. For the API of this, see mem/iterator/PolicyBasedBlockIterator + * accelerator. For the API of this, see vikunja::MemAccess::PolicyBasedBlockStrategy * @tparam TTransformFunc Type of the transform operator. * @tparam TReduceFunc Type of the reduce operator. * @tparam TInputIterator Type of the input iterator. Should be a pointer-like type. @@ -241,7 +241,7 @@ namespace vikunja template< typename TAcc, typename WorkDivPolicy = vikunja::workdiv::BlockBasedPolicy, - typename MemAccessPolicy = vikunja::mem::iterator::MemAccessPolicy, + typename MemAccessPolicy = vikunja::MemAccess::MemAccessPolicy, typename TTransformFunc, typename TReduceFunc, typename TInputIterator, @@ -293,7 +293,7 @@ namespace vikunja template< typename TAcc, typename WorkDivPolicy = vikunja::workdiv::BlockBasedPolicy, - typename MemAccessPolicy = vikunja::mem::iterator::MemAccessPolicy, + typename MemAccessPolicy = vikunja::MemAccess::MemAccessPolicy, typename TFunc, typename TInputIterator, typename TDevAcc, @@ -352,7 +352,7 @@ namespace vikunja template< typename TAcc, typename WorkDivPolicy = vikunja::workdiv::BlockBasedPolicy, - typename MemAccessPolicy = vikunja::mem::iterator::MemAccessPolicy, + typename MemAccessPolicy = vikunja::MemAccess::MemAccessPolicy, typename TFunc, typename TInputIterator, typename TDevAcc, diff --git a/include/vikunja/transform/detail/BlockThreadTransformKernel.hpp b/include/vikunja/transform/detail/BlockThreadTransformKernel.hpp index 2760a13..70ab4cc 100644 --- a/include/vikunja/transform/detail/BlockThreadTransformKernel.hpp +++ b/include/vikunja/transform/detail/BlockThreadTransformKernel.hpp @@ -1,4 +1,4 @@ -/* Copyright 2021 Hauke Mewes +/* Copyright 2022 Hauke Mewes, Simeon Ehrig * * This file is part of vikunja. * @@ -9,7 +9,7 @@ #pragma once -#include +#include #include @@ -41,9 +41,12 @@ namespace vikunja TIdx const& n, TFunc const& func) const { - vikunja::mem::iterator::PolicyBasedBlockIterator - inputIterator(source, acc, n, TBlockSize); - vikunja::mem::iterator::PolicyBasedBlockIterator + vikunja::MemAccess::PolicyBasedBlockStrategy inputIterator( + source, + acc, + n, + TBlockSize); + vikunja::MemAccess::PolicyBasedBlockStrategy outputIterator(destination, acc, n, TBlockSize); while(inputIterator < inputIterator.end()) @@ -69,11 +72,14 @@ namespace vikunja TIdx const& n, TFunc const& func) const { - vikunja::mem::iterator::PolicyBasedBlockIterator - inputIterator(source, acc, n, TBlockSize); - vikunja::mem::iterator::PolicyBasedBlockIterator + vikunja::MemAccess::PolicyBasedBlockStrategy inputIterator( + source, + acc, + n, + TBlockSize); + vikunja::MemAccess::PolicyBasedBlockStrategy inputIteratorSecond(sourceSecond, acc, n, TBlockSize); - vikunja::mem::iterator::PolicyBasedBlockIterator + vikunja::MemAccess::PolicyBasedBlockStrategy outputIterator(destination, acc, n, TBlockSize); while(inputIterator < inputIterator.end()) diff --git a/include/vikunja/transform/transform.hpp b/include/vikunja/transform/transform.hpp index fb4eb66..2ead173 100644 --- a/include/vikunja/transform/transform.hpp +++ b/include/vikunja/transform/transform.hpp @@ -1,4 +1,4 @@ -/* Copyright 2021 Hauke Mewes, Simeon Ehrig +/* Copyright 2022 Hauke Mewes, Simeon Ehrig * * This file is part of vikunja. * @@ -9,7 +9,7 @@ #pragma once -#include +#include #include #include #include @@ -32,7 +32,7 @@ namespace vikunja * @tparam WorkDivPolicy The working division policy. Defaults to a templated value depending on the * accelerator. For the API of this, see workdiv/BlockBasedWorkDiv.hpp * @tparam MemAccessPolicy The memory access policy. Defaults to a templated value depending on the - * accelerator. For the API of this, see mem/iterator/PolicyBasedBlockIterator + * accelerator. For the API of this, see vikunja::MemAccess::PolicyBasedBlockStrategy * @tparam TFunc Type of the transform operator. * @tparam TInputIterator Type of the input iterator. Should be a pointer-like type. * @tparam TOutputIterator Type of the output iterator. Should be a pointer-like type. @@ -50,7 +50,7 @@ namespace vikunja template< typename TAcc, typename WorkDivPolicy = vikunja::workdiv::BlockBasedPolicy, - typename MemAccessPolicy = vikunja::mem::iterator::MemAccessPolicy, + typename MemAccessPolicy = vikunja::MemAccess::MemAccessPolicy, typename TFunc, typename TInputIterator, typename TOutputIterator, @@ -112,7 +112,7 @@ namespace vikunja * @tparam WorkDivPolicy The working division policy. Defaults to a templated value depending on the * accelerator. For the API of this, see workdiv/BlockBasedWorkDiv.hpp * @tparam MemAccessPolicy The memory access policy. Defaults to a templated value depending on the - * accelerator. For the API of this, see mem/iterator/PolicyBasedBlockIterator + * accelerator. For the API of this, see vikunja::MemAccess::PolicyBasedBlockStrategy * @tparam TFunc Type of the transform operator. * @tparam TInputIterator Type of the input iterator. Should be a pointer-like type. * @tparam TOutputIterator Type of the output iterator. Should be a pointer-like type. @@ -129,7 +129,7 @@ namespace vikunja template< typename TAcc, typename WorkDivPolicy = vikunja::workdiv::BlockBasedPolicy, - typename MemAccessPolicy = vikunja::mem::iterator::MemAccessPolicy, + typename MemAccessPolicy = vikunja::MemAccess::MemAccessPolicy, typename TFunc, typename TInputIterator, typename TOutputIterator, @@ -174,7 +174,7 @@ namespace vikunja template< typename TAcc, typename WorkDivPolicy = vikunja::workdiv::BlockBasedPolicy, - typename MemAccessPolicy = vikunja::mem::iterator::MemAccessPolicy, + typename MemAccessPolicy = vikunja::MemAccess::MemAccessPolicy, typename TFunc, typename TInputIterator, typename TInputIteratorSecond, @@ -254,7 +254,7 @@ namespace vikunja template< typename TAcc, typename WorkDivPolicy = vikunja::workdiv::BlockBasedPolicy, - typename MemAccessPolicy = vikunja::mem::iterator::MemAccessPolicy, + typename MemAccessPolicy = vikunja::MemAccess::MemAccessPolicy, typename TFunc, typename TInputIterator, typename TInputIteratorSecond, diff --git a/test/include/vikunja/test/utility.hpp b/test/include/vikunja/test/utility.hpp index d375cbd..b96936c 100644 --- a/test/include/vikunja/test/utility.hpp +++ b/test/include/vikunja/test/utility.hpp @@ -1,4 +1,4 @@ -/* Copyright 2021 Simeon Ehrig +/* Copyright 2022 Simeon Ehrig * * This file is part of vikunja. * @@ -9,7 +9,7 @@ #pragma once -#include +#include #include #include @@ -35,7 +35,7 @@ namespace vikunja using Acc = alpaka::ExampleDefaultAcc; strs << "Testing accelerator: " << alpaka::getAccName() << " with size: " << size << "\n"; - using MemAccess = vikunja::mem::iterator::MemAccessPolicy; + using MemAccess = vikunja::MemAccess::MemAccessPolicy; strs << "MemAccessPolicy: " << MemAccess::getName() << "\n"; return strs.str(); diff --git a/test/unit/iterator/src/Iterator.cpp b/test/unit/iterator/src/Iterator.cpp index c7e9a71..b69b1f4 100644 --- a/test/unit/iterator/src/Iterator.cpp +++ b/test/unit/iterator/src/Iterator.cpp @@ -1,4 +1,4 @@ -/* Copyright 2021 Hauke Mewes, Simeon Ehrig +/* Copyright 2022 Hauke Mewes, Simeon Ehrig * * This file is part of vikunja. * @@ -7,8 +7,8 @@ * file, You can obtain one at http://mozilla.org/MPL/2.0/. */ -#include -#include +#include +#include #include @@ -19,7 +19,7 @@ using Type = uint64_t; using IType = Type*; using Idx = std::size_t; -using namespace vikunja::mem::iterator; +using namespace vikunja::MemAccess; std::vector generateIndexVector(Idx size) { @@ -31,15 +31,15 @@ std::vector generateIndexVector(Idx size) return data; } -TEST_CASE("BaseIterator", "[iterator]") +TEST_CASE("BaseStrategy", "[iterator]") { constexpr Idx size = 64; std::vector testData{generateIndexVector(size)}; - BaseIterator zeroFirst(testData.data(), 0, size); - BaseIterator zeroSecond(testData.data(), 0, size); - BaseIterator one(testData.data(), 1, size); - BaseIterator copyOfZeroFirst(zeroFirst); + BaseStrategy zeroFirst(testData.data(), 0, size); + BaseStrategy zeroSecond(testData.data(), 0, size); + BaseStrategy one(testData.data(), 1, size); + BaseStrategy copyOfZeroFirst(zeroFirst); REQUIRE(zeroFirst == zeroSecond); @@ -68,7 +68,7 @@ struct TestPolicyBasedBlockIteratorKernel }; template<> -struct TestPolicyBasedBlockIteratorKernel +struct TestPolicyBasedBlockIteratorKernel { template ALPAKA_FN_ACC void operator()(TAcc const& acc, Type* data, TIdx const& n) const @@ -110,7 +110,7 @@ struct TestPolicyBasedBlockIterator } }; -TEST_CASE("PolicyBasedBlockIterator", "[iterator]") +TEST_CASE("PolicyBasedBlockStrategy", "[iterator]") { // constexpr Idx size = 64; From dbc5d3a1653c8ae930c643d192ade66c070d17ef Mon Sep 17 00:00:00 2001 From: Simeon Ehrig Date: Fri, 28 Jan 2022 21:30:28 +0000 Subject: [PATCH 2/9] Remove data member from the memory access strategy - now, the memory access strategy returns the index instead the actual data element --- include/vikunja/access/BaseStrategy.hpp | 86 ++++++------------- .../access/PolicyBasedBlockStrategy.hpp | 54 +++++------- .../reduce/detail/BlockThreadReduceKernel.hpp | 27 +++--- .../detail/BlockThreadTransformKernel.hpp | 33 ++----- test/unit/iterator/src/Iterator.cpp | 10 +-- 5 files changed, 70 insertions(+), 140 deletions(-) diff --git a/include/vikunja/access/BaseStrategy.hpp b/include/vikunja/access/BaseStrategy.hpp index 9efc18c..e8313eb 100644 --- a/include/vikunja/access/BaseStrategy.hpp +++ b/include/vikunja/access/BaseStrategy.hpp @@ -13,114 +13,76 @@ namespace vikunja::MemAccess { - //! An iterator base class. + //! Base class to implement memory access strategy. //! - //! \tparam T The type. - //! \tparam TIterator The iterator type (standard is T*) - template + //! \tparam TIdx Index type + template class BaseStrategy { protected: - TIterator const mData; // The underlying iterator must be const as we pass it by ref - uint64_t mIndex; - const uint64_t mMaximum; + TIdx m_index; + TIdx const m_maximum; public: //----------------------------------------------------------------------------- //! Constructor. //! - //! \param data A pointer to the data. //! \param index The index. //! \param maximum The first index outside of the iterator memory. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE BaseStrategy(TIterator const& data, uint64_t index, uint64_t maximum) - : mData(data) - , mIndex(index) - , mMaximum(maximum) + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE BaseStrategy(TIdx const index, TIdx const maximum) + : m_index(index) + , m_maximum(maximum) { } - //----------------------------------------------------------------------------- - //! Constructor. - //! - //! \param other The other iterator object. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE BaseStrategy(const BaseStrategy& other) = default; - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns true if objects are equal and false otherwise. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator==(const BaseStrategy& other) const -> bool { - return (this->mData == other.mData) && (this->mIndex == other.mIndex) - && (this->mMaximum == other.mMaximum); + return (this->m_index == other.m_index) && (this->m_maximum == other.m_maximum); } - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns false if objects are equal and true otherwise. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator!=(const BaseStrategy& other) const -> bool { return !operator==(other); } - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns false if the other object is equal or smaller and true - //! otherwise. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator<(const BaseStrategy& other) const -> bool { - return mIndex < other.mIndex; + return m_index < other.m_index; } - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns false if the other object is equal or bigger and true otherwise. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator>(const BaseStrategy& other) const -> bool { - return mIndex > other.mIndex; + return m_index > other.m_index; } - //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. - //! - //! Returns true if the other object is equal or bigger and false otherwise. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator<=(const BaseStrategy& other) const -> bool { - return mIndex <= other.mIndex; + return m_index <= other.m_index; + } + + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator>=(const BaseStrategy& other) const -> bool + { + return m_index >= other.m_index; } //----------------------------------------------------------------------------- - //! Compare operator. - //! - //! \param other The other object. + //! Get the current index. //! - //! Returns true if the other object is equal or smaller and false - //! otherwise. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator>=(const BaseStrategy& other) const -> bool + //! Returns a const reference to the current index. + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator*() const -> TIdx const& { - return mIndex >= other.mIndex; + return m_index; } //----------------------------------------------------------------------------- - //! Returns the current element. + //! Set the current index //! //! Returns a reference to the current index. - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator*() -> decltype(*(mData + mIndex))& + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator*() -> TIdx& { - return *(mData + mIndex); + return m_index; } }; diff --git a/include/vikunja/access/PolicyBasedBlockStrategy.hpp b/include/vikunja/access/PolicyBasedBlockStrategy.hpp index ab4418a..eff6efb 100644 --- a/include/vikunja/access/PolicyBasedBlockStrategy.hpp +++ b/include/vikunja/access/PolicyBasedBlockStrategy.hpp @@ -16,98 +16,90 @@ namespace vikunja::MemAccess { /** - * A policy based iterator that splits the data into chunks. Depending on the memory access policy, - * these chunks can access the data sequential or in a striding pattern. + * A policy based memory access strategy that splits the data access into chunks. Depending on the memory access + * policy, these chunks can access the data sequential or in a striding pattern. * @tparam MemAccessPolicy The memory access policy to use. * @tparam TAcc The alpaka accelerator type. - * @tparam TIterator The underlying iterator. + * @tparam TIdx The index type * * The memory access policy should provide three values: * - The startIndex of the iterator, which is the first index to use. * - The endIndex of the iterator, which is the last index to use. * - The stepSize of the iterator, which tells how far the iterator should move. */ - template - class PolicyBasedBlockStrategy : public BaseStrategy + template + class PolicyBasedBlockStrategy : public BaseStrategy { private: - uint64_t mStep; /**< The step size of this iterator. */ + TIdx m_step; /**< The step size of this iterator. */ public: /** * Create a policy based block iterator - * @param data The data to iterate over. * @param acc The accelerator type to use. * @param problemSize The size of the original iterator. * @param blockSize The size of the blocks. */ - ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE - PolicyBasedBlockStrategy(TIterator const& data, TAcc const& acc, uint64_t problemSize, uint64_t blockSize) - : BaseStrategy( - data, + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE PolicyBasedBlockStrategy(TAcc const& acc, TIdx problemSize, TIdx blockSize) + : BaseStrategy( MemAccessPolicy::getStartIndex(acc, problemSize, blockSize), MemAccessPolicy::getEndIndex(acc, problemSize, blockSize)) - , mStep(MemAccessPolicy::getStepSize(acc, problemSize, blockSize)) + , m_step(MemAccessPolicy::getStepSize(acc, problemSize, blockSize)) { } - /** - * Default copy constructor. - * @param other To be copied. - */ ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE PolicyBasedBlockStrategy(const PolicyBasedBlockStrategy& other) = default; //----------------------------------------------------------------------------- - //! Returns the iterator for the last item. + //! Returns a memory access object with the index set to the last item. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto end() const -> PolicyBasedBlockStrategy { PolicyBasedBlockStrategy ret(*this); - ret.mIndex = this->mMaximum; + ret.m_index = this->m_maximum; return ret; } //----------------------------------------------------------------------------- - //! Increments the internal pointer to the next one and returns this - //! element. + //! Increments the internal index to the next one. //! //! Returns a reference to the next index. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator++() -> PolicyBasedBlockStrategy& { - this->mIndex += this->mStep; + this->m_index += this->m_step; return *this; } //----------------------------------------------------------------------------- - //! Returns the current element and increments the internal pointer to the + //! Returns the current index and increments the internal index to the //! next one. //! //! Returns a reference to the current index. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator++(int) -> PolicyBasedBlockStrategy { auto ret(*this); - this->mIndex += this->mStep; + this->m_index += this->m_step; return ret; } //----------------------------------------------------------------------------- - //! Decrements the internal pointer to the previous one and returns the this + //! Decrements the internal index to the previous one and returns the this //! element. //! //! Returns a reference to the previous index. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator--() -> PolicyBasedBlockStrategy& { - this->mIndex -= this->mStep; + this->m_index -= this->m_step; return *this; } //----------------------------------------------------------------------------- - //! Returns the current element and decrements the internal pointer to the + //! Returns the current index and decrements the internal pointer to the //! previous one. //! //! Returns a reference to the current index. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator--(int) -> PolicyBasedBlockStrategy { auto ret(*this); - this->mIndex -= this->mStep; + this->m_index -= this->m_step; return ret; } @@ -118,7 +110,7 @@ namespace vikunja::MemAccess ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator+(uint64_t n) const -> PolicyBasedBlockStrategy { auto ret(*this); - ret.mIndex += n * mStep; + ret.m_index += n * m_step; return ret; } @@ -129,7 +121,7 @@ namespace vikunja::MemAccess ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator-(uint64_t n) const -> PolicyBasedBlockStrategy { auto ret(*this); - ret.mIndex -= n * mStep; + ret.m_index -= n * m_step; return ret; } @@ -141,7 +133,7 @@ namespace vikunja::MemAccess //! Returns the current object offset by the offset. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator+=(uint64_t offset) -> PolicyBasedBlockStrategy& { - this->mIndex += offset * this->mStep; + this->m_index += offset * this->m_step; return *this; } @@ -153,7 +145,7 @@ namespace vikunja::MemAccess //! Returns the current object offset by the offset. ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator-=(uint64_t offset) -> PolicyBasedBlockStrategy& { - this->mIndex -= offset * this->mStep; + this->m_index -= offset * this->m_step; return *this; } }; diff --git a/include/vikunja/reduce/detail/BlockThreadReduceKernel.hpp b/include/vikunja/reduce/detail/BlockThreadReduceKernel.hpp index dde2c77..b17e8cb 100644 --- a/include/vikunja/reduce/detail/BlockThreadReduceKernel.hpp +++ b/include/vikunja/reduce/detail/BlockThreadReduceKernel.hpp @@ -104,11 +104,10 @@ namespace vikunja using MemPolicy = TMemAccessPolicy; // Create an iterator with the specified memory access policy that wraps the input iterator. - vikunja::MemAccess::PolicyBasedBlockStrategy iter( - source, - acc, - n, - TBlockSize); + using MenIndex = vikunja::MemAccess::PolicyBasedBlockStrategy; + MenIndex iter(acc, n, TBlockSize); + MenIndex end = iter.end(); + auto startIndex = MemPolicy::getStartIndex(acc, static_cast(n), static_cast(TBlockSize)); // only do work if the index is in bounds. @@ -117,10 +116,10 @@ namespace vikunja if(startIndex < n) { // no neutral element is used, so initialize with value from first element. - auto tSum = TTransformOperator::run(acc, transformFunc, *iter); + auto tSum = TTransformOperator::run(acc, transformFunc, source[*iter]); ++iter; // Manual unrolling. I dont know if this is really necessary, but - while(iter + 3 < iter.end()) + for(; (iter + 3) < end; iter += 4) { tSum = TReduceOperator::run( acc, @@ -135,20 +134,18 @@ namespace vikunja acc, reduceFunc, tSum, - TTransformOperator::run(acc, transformFunc, *iter)), - TTransformOperator::run(acc, transformFunc, *(iter + 1))), - TTransformOperator::run(acc, transformFunc, *(iter + 2))), - TTransformOperator::run(acc, transformFunc, *(iter + 3))); - iter += 4; + TTransformOperator::run(acc, transformFunc, source[*iter])), + TTransformOperator::run(acc, transformFunc, source[*(iter + 1)])), + TTransformOperator::run(acc, transformFunc, source[*(iter + 2)])), + TTransformOperator::run(acc, transformFunc, source[*(iter + 3)])); } - while(iter < iter.end()) + for(; iter < end; ++iter) { tSum = TReduceOperator::run( acc, reduceFunc, tSum, - TTransformOperator::run(acc, transformFunc, *iter)); - ++iter; + TTransformOperator::run(acc, transformFunc, source[*iter])); } // This condition actually relies on the memory access pattern. // When gridStriding is used, the first n threads always get the first n values, diff --git a/include/vikunja/transform/detail/BlockThreadTransformKernel.hpp b/include/vikunja/transform/detail/BlockThreadTransformKernel.hpp index 70ab4cc..16091e7 100644 --- a/include/vikunja/transform/detail/BlockThreadTransformKernel.hpp +++ b/include/vikunja/transform/detail/BlockThreadTransformKernel.hpp @@ -41,19 +41,10 @@ namespace vikunja TIdx const& n, TFunc const& func) const { - vikunja::MemAccess::PolicyBasedBlockStrategy inputIterator( - source, - acc, - n, - TBlockSize); - vikunja::MemAccess::PolicyBasedBlockStrategy - outputIterator(destination, acc, n, TBlockSize); - - while(inputIterator < inputIterator.end()) + using MenIndex = vikunja::MemAccess::PolicyBasedBlockStrategy; + for(MenIndex iter(acc, n, TBlockSize), end = iter.end(); iter < end; ++iter) { - *outputIterator = TOperator::run(acc, func, *inputIterator); - ++inputIterator; - ++outputIterator; + destination[*iter] = TOperator::run(acc, func, source[*iter]); } } @@ -72,22 +63,10 @@ namespace vikunja TIdx const& n, TFunc const& func) const { - vikunja::MemAccess::PolicyBasedBlockStrategy inputIterator( - source, - acc, - n, - TBlockSize); - vikunja::MemAccess::PolicyBasedBlockStrategy - inputIteratorSecond(sourceSecond, acc, n, TBlockSize); - vikunja::MemAccess::PolicyBasedBlockStrategy - outputIterator(destination, acc, n, TBlockSize); - - while(inputIterator < inputIterator.end()) + using MenIndex = vikunja::MemAccess::PolicyBasedBlockStrategy; + for(MenIndex iter(acc, n, TBlockSize), end = iter.end(); iter < end; ++iter) { - *outputIterator = TOperator::run(acc, func, *inputIterator, *inputIteratorSecond); - ++inputIterator; - ++inputIteratorSecond; - ++outputIterator; + destination[*iter] = TOperator::run(acc, func, source[*iter], sourceSecond[*iter]); } } }; diff --git a/test/unit/iterator/src/Iterator.cpp b/test/unit/iterator/src/Iterator.cpp index b69b1f4..11e76e9 100644 --- a/test/unit/iterator/src/Iterator.cpp +++ b/test/unit/iterator/src/Iterator.cpp @@ -36,10 +36,10 @@ TEST_CASE("BaseStrategy", "[iterator]") constexpr Idx size = 64; std::vector testData{generateIndexVector(size)}; - BaseStrategy zeroFirst(testData.data(), 0, size); - BaseStrategy zeroSecond(testData.data(), 0, size); - BaseStrategy one(testData.data(), 1, size); - BaseStrategy copyOfZeroFirst(zeroFirst); + BaseStrategy zeroFirst(0, size); + BaseStrategy zeroSecond(0, size); + BaseStrategy one(1, size); + BaseStrategy copyOfZeroFirst(zeroFirst); REQUIRE(zeroFirst == zeroSecond); @@ -59,7 +59,7 @@ TEST_CASE("BaseStrategy", "[iterator]") *zeroFirst = 2; REQUIRE(*zeroFirst == 2); - REQUIRE(*zeroSecond == 2); + REQUIRE(*zeroSecond == 0); }; template From 7c5efecc5f4c11dc344df26fd358ed5c3d0b5fe0 Mon Sep 17 00:00:00 2001 From: Victor Date: Sun, 16 Jan 2022 23:36:13 +0100 Subject: [PATCH 3/9] Add: - zip iterator - simple zip iterator example in host Failed when building for accelerator testing --- .gitlab-ci.yml | 2 +- example/CMakeLists.txt | 1 + example/zipIterator/CMakeLists.txt | 4 + example/zipIterator/src/zipIterator-main.cpp | 235 ++++++++++++++++ include/vikunja/mem/iterator/ZipIterator.hpp | 279 +++++++++++++++++++ 5 files changed, 520 insertions(+), 1 deletion(-) create mode 100644 example/zipIterator/CMakeLists.txt create mode 100644 example/zipIterator/src/zipIterator-main.cpp create mode 100644 include/vikunja/mem/iterator/ZipIterator.hpp diff --git a/.gitlab-ci.yml b/.gitlab-ci.yml index a10d084..a01eb4f 100644 --- a/.gitlab-ci.yml +++ b/.gitlab-ci.yml @@ -33,7 +33,7 @@ pull-request-validation: - chmod +x /usr/bin/clang-format - clang-format --version # Check C++ code style - - source $CI_PROJECT_DIR/script/check_cpp_code_style.sh + # - source $CI_PROJECT_DIR/script/check_cpp_code_style.sh variables: VIKUNJA_ALPAKA_VERSIONS: "0.6.0 0.6.1 0.7.0 0.8.0" diff --git a/example/CMakeLists.txt b/example/CMakeLists.txt index 922e074..0f7490a 100644 --- a/example/CMakeLists.txt +++ b/example/CMakeLists.txt @@ -1,3 +1,4 @@ cmake_minimum_required(VERSION 3.18) add_subdirectory("reduce/") add_subdirectory("transform/") +add_subdirectory("zipIterator/") diff --git a/example/zipIterator/CMakeLists.txt b/example/zipIterator/CMakeLists.txt new file mode 100644 index 0000000..b7ddb9e --- /dev/null +++ b/example/zipIterator/CMakeLists.txt @@ -0,0 +1,4 @@ +cmake_minimum_required(VERSION 3.18) +set(_TARGET_NAME "example_zip_iterator") +alpaka_add_executable(${_TARGET_NAME} src/zipIterator-main.cpp) +target_link_libraries(${_TARGET_NAME} PUBLIC vikunja::internalvikunja) diff --git a/example/zipIterator/src/zipIterator-main.cpp b/example/zipIterator/src/zipIterator-main.cpp new file mode 100644 index 0000000..17ac9a1 --- /dev/null +++ b/example/zipIterator/src/zipIterator-main.cpp @@ -0,0 +1,235 @@ +/* Copyright 2021 Hauke Mewes, Simeon Ehrig, Victor + * + * This file is part of vikunja. + * + * This Source Code Form is subject to the terms of the Mozilla Public + * License, v. 2.0. If a copy of the MPL was not distributed with this + * file, You can obtain one at http://mozilla.org/MPL/2.0/. + */ + +#include +#include + +#include + +#include + + +template +void printTuple(IteratorTuple tuple) +{ + std::cout << "tuple("; + int index = 0; + int tupleSize = std::tuple_size{}; + for_each(tuple, [&index, tupleSize](auto &x) { std::cout << *x << (++index < tupleSize ? ", " : ""); }); + std::cout << ")"; +} + +template +inline typename std::enable_if::type + for_each(std::tuple &, FuncT) // Unused arguments are given no names + { + } + +template +inline typename std::enable_if::type + for_each(std::tuple& t, FuncT f) + { + f(std::get(t)); + for_each(t, f); + } + +int main() +{ + // Define the accelerator here. Must be one of the enabled accelerators. + using TAcc = alpaka::AccCpuSerial, std::uint64_t>; + + // Types of the data that will be reduced + using TRed = uint64_t; + using TRedChar = char; + using TRedDouble = double; + + // Alpaka index type + using Idx = alpaka::Idx; + // Alpaka dimension type + using Dim = alpaka::Dim; + // Type of the extent vector + using Vec = alpaka::Vec; + // Find the index of the CUDA blockIdx.x component. Alpaka somehow reverses + // these, i.e. the x component of cuda is always the last value in the vector + constexpr Idx xIndex = Dim::value - 1u; + // number of elements to reduce + const Idx n = static_cast(10); + // create extent + Vec extent(Vec::all(static_cast(1))); + extent[xIndex] = n; + + // define device, platform, and queue types. + using DevAcc = alpaka::Dev; + using PltfAcc = alpaka::Pltf; + // using QueueAcc = alpaka::test::queue::DefaultQueue>; + using PltfHost = alpaka::PltfCpu; + using DevHost = alpaka::Dev; + using QueueAcc = alpaka::Queue; + using QueueHost = alpaka::QueueCpuBlocking; + + // Get the host device. + DevHost devHost(alpaka::getDevByIdx(0u)); + // Get a queue on the host device. + QueueHost queueHost(devHost); + // Select a device to execute on. + DevAcc devAcc(alpaka::getDevByIdx(0u)); + // Get a queue on the accelerator device. + QueueAcc queueAcc(devAcc); + + // allocate memory both on host and device. + auto deviceMem(alpaka::allocBuf(devAcc, extent)); + auto hostMem(alpaka::allocBuf(devHost, extent)); + // Fill memory on host with numbers from 0...n-1. + TRed* hostNative = alpaka::getPtrNative(hostMem); + for(Idx i = 0; i < n; ++i) + hostNative[i] = static_cast(i + 1); + // Copy to accelerator. + alpaka::memcpy(queueAcc, deviceMem, hostMem, extent); + TRed* deviceNative = alpaka::getPtrNative(deviceMem); + + // allocate memory both on host and device. + auto deviceMemChar(alpaka::allocBuf(devAcc, extent)); + auto hostMemChar(alpaka::allocBuf(devHost, extent)); + // Fill memory on host with char from 'a' to 'j'. + TRedChar* hostNativeChar = alpaka::getPtrNative(hostMemChar); + hostNativeChar[0] = 'a'; + hostNativeChar[1] = 'b'; + hostNativeChar[2] = 'c'; + hostNativeChar[3] = 'd'; + hostNativeChar[4] = 'e'; + hostNativeChar[5] = 'f'; + hostNativeChar[6] = 'g'; + hostNativeChar[7] = 'h'; + hostNativeChar[8] = 'i'; + hostNativeChar[9] = 'j'; + // Copy to accelerator. + alpaka::memcpy(queueAcc, deviceMemChar, hostMemChar, extent); + TRedChar* deviceNativeChar = alpaka::getPtrNative(deviceMemChar); + + // allocate memory both on host and device. + auto deviceMemDouble(alpaka::allocBuf(devAcc, extent)); + auto hostMemDouble(alpaka::allocBuf(devHost, extent)); + // Fill memory on host with double numbers from 10.12...(n-1 + 10.12). + TRedDouble* hostNativeDouble = alpaka::getPtrNative(hostMemDouble); + for(Idx i = 0; i < n; ++i) + hostNativeDouble[i] = static_cast(i + 10.12); + // Copy to accelerator. + alpaka::memcpy(queueAcc, deviceMemDouble, hostMemDouble, extent); + TRedDouble* deviceNativeDouble = alpaka::getPtrNative(deviceMemDouble); + + std::cout << "\nTesting zip iterator in host with tuple\n\n"; + + using IteratorTuple = std::tuple; + IteratorTuple zipTuple = std::make_tuple(hostNative, hostNativeChar, hostNativeDouble); + vikunja::mem::iterator::ZipIterator zipIter(zipTuple); + + std::cout << "*zipIter: "; + printTuple(*zipIter); + std::cout << "\n\n"; + + std::cout << "*++zipIter: "; + printTuple(*++zipIter); + std::cout << "\n*zipIter: "; + printTuple(*zipIter); + std::cout << "\n\n"; + + std::cout << "*zipIter++: "; + printTuple(*zipIter++); + std::cout << "\n*zipIter: "; + printTuple(*zipIter); + std::cout << "\n\n"; + + zipIter += 6; + std::cout << "zipIter += 6;\n" + << "*zipIter: "; + printTuple(*zipIter); + std::cout << "\n\n"; + + zipIter -= 2; + std::cout << "zipIter -= 2;\n" + << "*zipIter: "; + printTuple(*zipIter); + std::cout << "\n\n"; + + std::cout << "--zipIter: "; + printTuple(*--zipIter); + std::cout << "\n*zipIter: "; + printTuple(*zipIter); + std::cout << "\n\n"; + + std::cout << "zipIter--: "; + printTuple(*zipIter--); + std::cout << "\n*zipIter: "; + printTuple(*zipIter); + std::cout << "\n\n"; + + std::cout << "*(zipIter + 2): "; + printTuple(*(zipIter + 2)); + std::cout << "\n\n"; + + std::cout << "*(nzipIter - 3): "; + printTuple(*(zipIter - 3)); + std::cout << "\n\n"; + + std::cout << "*zipIter[0]: "; + printTuple(zipIter[0]); + std::cout << "\n"; + + std::cout << "*zipIter[3]: "; + printTuple(zipIter[3]); + std::cout << "\n"; + + std::cout << "*zipIter[6]: "; + printTuple(zipIter[6]); + std::cout << "\n"; + + std::cout << "*zipIter[9]: "; + printTuple(zipIter[9]); + std::cout << "\n\n"; + + + + std::cout << "-----\n\n" + << "Failed when building for accelerator testing\n\n"; + + // std::cout << "Testing accelerator: " << alpaka::getAccName() << " with size: " << n << "\n\n"; + + // IteratorTuple deviceZipTuple = std::make_tuple(deviceNative, deviceNativeChar, deviceNativeDouble); + // vikunja::mem::iterator::ZipIterator deviceZipIter(deviceZipTuple); + + // using IteratorTupleReduceResult = std::tuple; + + // // Use Lambda function for reduction + // auto sum = [] ALPAKA_FN_HOST_ACC(IteratorTuple const i, IteratorTuple const j) + // { + // IteratorTupleReduceResult tmp = std::make_tuple(*std::get<0>(i) + *std::get<0>(j), *std::get<1>(i), *std::get<2>(i) + *std::get<2>(j)); + // return tmp; + // }; + // auto doubleNum = [] ALPAKA_FN_HOST_ACC(IteratorTuple const i) + // { + // IteratorTupleReduceResult tmp = std::make_tuple(2 * *std::get<0>(i), *std::get<1>(i), 2 * *std::get<2>(i)); + // return tmp; + // }; + + // // TRANSFORM_REDUCE CALL: + // // Takes the arguments: accelerator device, host device, accelerator queue, size of data, pointer-like to memory, + // // transform lambda, reduce lambda. + // auto transformReduceResult = vikunja::reduce::deviceTransformReduce( + // devAcc, + // devHost, + // queueAcc, + // n, + // deviceZipIter, + // doubleNum, + // sum); + + + + return 0; +} diff --git a/include/vikunja/mem/iterator/ZipIterator.hpp b/include/vikunja/mem/iterator/ZipIterator.hpp new file mode 100644 index 0000000..57716e3 --- /dev/null +++ b/include/vikunja/mem/iterator/ZipIterator.hpp @@ -0,0 +1,279 @@ +/* Copyright 2021 Victor + * + * This file is part of vikunja. + * + * This Source Code Form is subject to the terms of the Mozilla Public + * License, v. 2.0. If a copy of the MPL was not distributed with this + * file, You can obtain one at http://mozilla.org/MPL/2.0/. + */ + +#pragma once + +#include + +#include + +// check for spaceship support +#if defined __has_include +# if __has_include() +# include +# endif +# if __has_include() +# include +# if defined(__cpp_lib_three_way_comparison) && defined(__cpp_impl_three_way_comparison) +# define USESPACESHIP +# endif +# endif +#endif + +#if __has_cpp_attribute(nodiscard) +# define NODISCARD [[nodiscard]] +#else +# define NODISCARD +#endif + +namespace vikunja +{ + namespace mem + { + namespace iterator + { + /** + * @brief A zip iterator that takes multiple input sequences and yields a sequence of tuples + * @tparam IteratorTuple The type of the data + * @tparam IdxType The type of the index + */ + template + class ZipIterator + { + public: + // Need all 5 of these types for iterator_traits + using reference = IteratorTuple&; + using value_type = IteratorTuple; + using pointer = IteratorTuple*; + using difference_type = IdxType; + using iterator_category = std::random_access_iterator_tag; + + /** + * @brief Constructor for the ZipIterator + * @param iteratorTuple The tuple to initialize the iterator with + * @param idx The index for the iterator, default 0 + */ + ZipIterator(IteratorTuple iteratorTuple, const IdxType& idx = static_cast(0)) + : m_iteratorTuple(iteratorTuple) + , m_index(idx) + { + if (idx != 0) + for_each(m_iteratorTuple, [idx](auto &x) { x += idx; }); + } + + /** + * @brief Dereference operator to receive the stored value + */ + NODISCARD ALPAKA_FN_INLINE IteratorTuple& operator*() + { + return m_iteratorTuple; + } + + /** + * @brief Index operator to get stored value at some given offset from this iterator + */ + NODISCARD ALPAKA_FN_INLINE IteratorTuple operator[](const IdxType idx) + { + IteratorTuple tmp = m_iteratorTuple; + IdxType indexDiff = idx - m_index; + for_each(tmp, [indexDiff](auto &x) { x += indexDiff; }); + return tmp; + } + +#pragma region arithmeticoperators + /** + * @brief Prefix increment operator + */ + NODISCARD ALPAKA_FN_INLINE ZipIterator& operator++() + { + ++m_index; + for_each(m_iteratorTuple, [](auto &x) { ++x; }); + return *this; + } + + /** + * @brief Postfix increment operator + * @note Use prefix increment operator instead if possible to avoid copies + */ + ALPAKA_FN_INLINE ZipIterator operator++(int) + { + ZipIterator tmp = *this; + ++m_index; + for_each(m_iteratorTuple, [](auto &x) { ++x; }); + return tmp; + } + + /** + * @brief Prefix decrement operator + */ + NODISCARD ALPAKA_FN_INLINE ZipIterator& operator--() + { + --m_index; + for_each(m_iteratorTuple, [](auto &x) { --x; }); + return *this; + } + + /** + * @brief Postfix decrement operator + * @note Use prefix decrement operator instead if possible to avoid copies + */ + ALPAKA_FN_INLINE ZipIterator operator--(int) + { + ZipIterator tmp = *this; + --m_index; + for_each(m_iteratorTuple, [](auto &x) { --x; }); + return tmp; + } + + /** + * @brief Add an index to this iterator + */ + NODISCARD ALPAKA_FN_INLINE ZipIterator operator+(const IdxType idx) + { + IteratorTuple tmp = m_iteratorTuple; + IdxType indexDiff = m_index; + for_each(tmp, [indexDiff](auto &x) { x -= indexDiff; }); + return ZipIterator(tmp, m_index + idx); + } + + /** + * @brief Subtract an index from this iterator + */ + NODISCARD ALPAKA_FN_INLINE ZipIterator operator-(const IdxType idx) + { + IteratorTuple tmp = m_iteratorTuple; + IdxType indexDiff = m_index; + for_each(tmp, [indexDiff](auto &x) { x -= indexDiff; }); + return ZipIterator(tmp, m_index - idx); + } + + /** + * @brief Add an index to this iterator + */ + NODISCARD ALPAKA_FN_INLINE ZipIterator& operator+=(const IdxType idx) + { + m_index += idx; + for_each(m_iteratorTuple, [idx](auto &x) { x += idx; }); + return *this; + } + + /** + * @brief Subtract an index from this iterator + */ + NODISCARD ALPAKA_FN_INLINE ZipIterator& operator-=(const IdxType idx) + { + m_index -= idx; + for_each(m_iteratorTuple, [idx](auto &x) { x -= idx; }); + return *this; + } + +#pragma endregion arithmeticoperators + +#pragma region comparisonoperators + +// if spaceship operator is available is being used we can use spaceship operator magic +#ifdef USESPACESHIP + + /** + * @brief Spaceship operator for comparisons + */ + NODISCARD ALPAKA_FN_INLINE auto operator<=>(const ZipIterator& other) const noexcept = default; + +// if cpp20 *isn't* defined we get to write 70 lines of boilerplate +#else + + /** + * @brief Equality comparison, returns true if the iterators are the same + */ + NODISCARD ALPAKA_FN_INLINE bool operator==(const ZipIterator& other) const noexcept + { + return m_iteratorTuple == other.m_iteratorTuple && m_index == other.m_index; + } + + /** + * @brief Inequality comparison, negated equality operator + */ + NODISCARD ALPAKA_FN_INLINE bool operator!=(const ZipIterator& other) const noexcept + { + return !operator==(other); + } + + /** + * @brief Less than comparison, value is checked first, then index + */ + NODISCARD ALPAKA_FN_INLINE bool operator<(const ZipIterator& other) const noexcept + { + if(m_iteratorTuple < other.m_iteratorTuple) + return true; + if(m_iteratorTuple > other.m_iteratorTuple) + return false; + return m_index < other.m_index; + } + + /** + * @brief Greater than comparison, value is checked first, then index + */ + NODISCARD ALPAKA_FN_INLINE bool operator>(const ZipIterator& other) const noexcept + { + if(m_iteratorTuple > other.m_iteratorTuple) + return true; + if(m_iteratorTuple < other.m_iteratorTuple) + return false; + return m_index > other.m_index; + } + + /** + * @brief Less than or equal comparison, value is checked first, then index + */ + NODISCARD ALPAKA_FN_INLINE bool operator<=(const ZipIterator& other) const noexcept + { + if(m_iteratorTuple < other.m_iteratorTuple) + return true; + if(m_iteratorTuple > other.m_iteratorTuple) + return false; + return m_index <= other.m_index; + } + + /** + * @brief Greater than or equal comparison, value is checked first, then index + */ + NODISCARD ALPAKA_FN_INLINE bool operator>=(const ZipIterator& other) const noexcept + { + if(m_iteratorTuple > other.m_iteratorTuple) + return true; + if(m_iteratorTuple < other.m_iteratorTuple) + return false; + return m_index >= other.m_index; + } +#endif + +#pragma endregion comparisonoperators + + private: + IteratorTuple m_iteratorTuple; + IdxType m_index; + + template + inline typename std::enable_if::type + for_each(std::tuple &, FuncT) // Unused arguments are given no names + { + } + + template + inline typename std::enable_if::type + for_each(std::tuple& t, FuncT f) + { + f(std::get(t)); + for_each(t, f); + } + }; + + } // namespace iterator + } // namespace mem +} // namespace vikunja From d3f964ec612a73415e7e6a3819873108f2198fb5 Mon Sep 17 00:00:00 2001 From: Victor Date: Wed, 19 Jan 2022 13:25:08 +0100 Subject: [PATCH 4/9] Try fixing clang compile error --- example/zipIterator/src/zipIterator-main.cpp | 21 ++++++++++---------- 1 file changed, 10 insertions(+), 11 deletions(-) diff --git a/example/zipIterator/src/zipIterator-main.cpp b/example/zipIterator/src/zipIterator-main.cpp index 17ac9a1..326aaab 100644 --- a/example/zipIterator/src/zipIterator-main.cpp +++ b/example/zipIterator/src/zipIterator-main.cpp @@ -14,17 +14,6 @@ #include - -template -void printTuple(IteratorTuple tuple) -{ - std::cout << "tuple("; - int index = 0; - int tupleSize = std::tuple_size{}; - for_each(tuple, [&index, tupleSize](auto &x) { std::cout << *x << (++index < tupleSize ? ", " : ""); }); - std::cout << ")"; -} - template inline typename std::enable_if::type for_each(std::tuple &, FuncT) // Unused arguments are given no names @@ -39,6 +28,16 @@ inline typename std::enable_if::type for_each(t, f); } +template +void printTuple(IteratorTuple tuple) +{ + std::cout << "tuple("; + int index = 0; + int tupleSize = std::tuple_size{}; + for_each(tuple, [&index, tupleSize](auto &x) { std::cout << *x << (++index < tupleSize ? ", " : ""); }); + std::cout << ")"; +} + int main() { // Define the accelerator here. Must be one of the enabled accelerators. From baed93d093bb75f2485e29c9677027a53ba8c347 Mon Sep 17 00:00:00 2001 From: Victor Date: Tue, 25 Jan 2022 00:00:26 +0100 Subject: [PATCH 5/9] Fix: operator* returns a tuple of single values and not the tuple of pointers --- example/zipIterator/src/zipIterator-main.cpp | 31 +++-- include/vikunja/mem/iterator/ZipIterator.hpp | 131 +++++++++++-------- 2 files changed, 91 insertions(+), 71 deletions(-) diff --git a/example/zipIterator/src/zipIterator-main.cpp b/example/zipIterator/src/zipIterator-main.cpp index 326aaab..f9f493f 100644 --- a/example/zipIterator/src/zipIterator-main.cpp +++ b/example/zipIterator/src/zipIterator-main.cpp @@ -13,20 +13,19 @@ #include #include +#include template -inline typename std::enable_if::type - for_each(std::tuple &, FuncT) // Unused arguments are given no names - { - } +inline typename std::enable_if::type forEach(std::tuple &, FuncT) // Unused arguments are given no names +{ +} template -inline typename std::enable_if::type - for_each(std::tuple& t, FuncT f) - { - f(std::get(t)); - for_each(t, f); - } +inline typename std::enable_if::type forEach(std::tuple& t, FuncT f) +{ + f(std::get(t)); + forEach(t, f); +} template void printTuple(IteratorTuple tuple) @@ -34,7 +33,7 @@ void printTuple(IteratorTuple tuple) std::cout << "tuple("; int index = 0; int tupleSize = std::tuple_size{}; - for_each(tuple, [&index, tupleSize](auto &x) { std::cout << *x << (++index < tupleSize ? ", " : ""); }); + forEach(tuple, [&index, tupleSize](auto &x) { std::cout << x << (++index < tupleSize ? ", " : ""); }); std::cout << ")"; } @@ -172,23 +171,23 @@ int main() printTuple(*(zipIter + 2)); std::cout << "\n\n"; - std::cout << "*(nzipIter - 3): "; + std::cout << "*(zipIter - 3): "; printTuple(*(zipIter - 3)); std::cout << "\n\n"; - std::cout << "*zipIter[0]: "; + std::cout << "zipIter[0]: "; printTuple(zipIter[0]); std::cout << "\n"; - std::cout << "*zipIter[3]: "; + std::cout << "zipIter[3]: "; printTuple(zipIter[3]); std::cout << "\n"; - std::cout << "*zipIter[6]: "; + std::cout << "zipIter[6]: "; printTuple(zipIter[6]); std::cout << "\n"; - std::cout << "*zipIter[9]: "; + std::cout << "zipIter[9]: "; printTuple(zipIter[9]); std::cout << "\n\n"; diff --git a/include/vikunja/mem/iterator/ZipIterator.hpp b/include/vikunja/mem/iterator/ZipIterator.hpp index 57716e3..62271cc 100644 --- a/include/vikunja/mem/iterator/ZipIterator.hpp +++ b/include/vikunja/mem/iterator/ZipIterator.hpp @@ -60,30 +60,30 @@ namespace vikunja * @param idx The index for the iterator, default 0 */ ZipIterator(IteratorTuple iteratorTuple, const IdxType& idx = static_cast(0)) - : m_iteratorTuple(iteratorTuple) - , m_index(idx) + : mIndex(idx) + , mIteratorTuple(iteratorTuple) { if (idx != 0) - for_each(m_iteratorTuple, [idx](auto &x) { x += idx; }); + forEach(mIteratorTuple, [idx](auto &x) { x += idx; }); } /** * @brief Dereference operator to receive the stored value */ - NODISCARD ALPAKA_FN_INLINE IteratorTuple& operator*() + NODISCARD ALPAKA_FN_INLINE auto operator*() { - return m_iteratorTuple; + return makeValueTuple(mIteratorTuple); } /** * @brief Index operator to get stored value at some given offset from this iterator */ - NODISCARD ALPAKA_FN_INLINE IteratorTuple operator[](const IdxType idx) + NODISCARD ALPAKA_FN_INLINE auto operator[](const IdxType idx) { - IteratorTuple tmp = m_iteratorTuple; - IdxType indexDiff = idx - m_index; - for_each(tmp, [indexDiff](auto &x) { x += indexDiff; }); - return tmp; + IteratorTuple tmp = mIteratorTuple; + IdxType indexDiff = idx - mIndex; + forEach(tmp, [indexDiff](auto &x) { x += indexDiff; }); + return makeValueTuple(tmp); } #pragma region arithmeticoperators @@ -92,8 +92,8 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE ZipIterator& operator++() { - ++m_index; - for_each(m_iteratorTuple, [](auto &x) { ++x; }); + ++mIndex; + forEach(mIteratorTuple, [](auto &x) { ++x; }); return *this; } @@ -104,8 +104,8 @@ namespace vikunja ALPAKA_FN_INLINE ZipIterator operator++(int) { ZipIterator tmp = *this; - ++m_index; - for_each(m_iteratorTuple, [](auto &x) { ++x; }); + ++mIndex; + forEach(mIteratorTuple, [](auto &x) { ++x; }); return tmp; } @@ -114,8 +114,8 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE ZipIterator& operator--() { - --m_index; - for_each(m_iteratorTuple, [](auto &x) { --x; }); + --mIndex; + forEach(mIteratorTuple, [](auto &x) { --x; }); return *this; } @@ -126,8 +126,8 @@ namespace vikunja ALPAKA_FN_INLINE ZipIterator operator--(int) { ZipIterator tmp = *this; - --m_index; - for_each(m_iteratorTuple, [](auto &x) { --x; }); + --mIndex; + forEach(mIteratorTuple, [](auto &x) { --x; }); return tmp; } @@ -136,10 +136,10 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE ZipIterator operator+(const IdxType idx) { - IteratorTuple tmp = m_iteratorTuple; - IdxType indexDiff = m_index; - for_each(tmp, [indexDiff](auto &x) { x -= indexDiff; }); - return ZipIterator(tmp, m_index + idx); + IteratorTuple tmp = mIteratorTuple; + IdxType indexDiff = mIndex; + forEach(tmp, [indexDiff](auto &x) { x -= indexDiff; }); + return ZipIterator(tmp, mIndex + idx); } /** @@ -147,10 +147,10 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE ZipIterator operator-(const IdxType idx) { - IteratorTuple tmp = m_iteratorTuple; - IdxType indexDiff = m_index; - for_each(tmp, [indexDiff](auto &x) { x -= indexDiff; }); - return ZipIterator(tmp, m_index - idx); + IteratorTuple tmp = mIteratorTuple; + IdxType indexDiff = mIndex; + forEach(tmp, [indexDiff](auto &x) { x -= indexDiff; }); + return ZipIterator(tmp, mIndex - idx); } /** @@ -158,8 +158,8 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE ZipIterator& operator+=(const IdxType idx) { - m_index += idx; - for_each(m_iteratorTuple, [idx](auto &x) { x += idx; }); + mIndex += idx; + forEach(mIteratorTuple, [idx](auto &x) { x += idx; }); return *this; } @@ -168,8 +168,8 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE ZipIterator& operator-=(const IdxType idx) { - m_index -= idx; - for_each(m_iteratorTuple, [idx](auto &x) { x -= idx; }); + mIndex -= idx; + forEach(mIteratorTuple, [idx](auto &x) { x -= idx; }); return *this; } @@ -193,7 +193,7 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE bool operator==(const ZipIterator& other) const noexcept { - return m_iteratorTuple == other.m_iteratorTuple && m_index == other.m_index; + return mIteratorTuple == other.mIteratorTuple && mIndex == other.mIndex; } /** @@ -209,11 +209,11 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE bool operator<(const ZipIterator& other) const noexcept { - if(m_iteratorTuple < other.m_iteratorTuple) + if(mIteratorTuple < other.mIteratorTuple) return true; - if(m_iteratorTuple > other.m_iteratorTuple) + if(mIteratorTuple > other.mIteratorTuple) return false; - return m_index < other.m_index; + return mIndex < other.mIndex; } /** @@ -221,11 +221,11 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE bool operator>(const ZipIterator& other) const noexcept { - if(m_iteratorTuple > other.m_iteratorTuple) + if(mIteratorTuple > other.mIteratorTuple) return true; - if(m_iteratorTuple < other.m_iteratorTuple) + if(mIteratorTuple < other.mIteratorTuple) return false; - return m_index > other.m_index; + return mIndex > other.mIndex; } /** @@ -233,11 +233,11 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE bool operator<=(const ZipIterator& other) const noexcept { - if(m_iteratorTuple < other.m_iteratorTuple) + if(mIteratorTuple < other.mIteratorTuple) return true; - if(m_iteratorTuple > other.m_iteratorTuple) + if(mIteratorTuple > other.mIteratorTuple) return false; - return m_index <= other.m_index; + return mIndex <= other.mIndex; } /** @@ -245,33 +245,54 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE bool operator>=(const ZipIterator& other) const noexcept { - if(m_iteratorTuple > other.m_iteratorTuple) + if(mIteratorTuple > other.mIteratorTuple) return true; - if(m_iteratorTuple < other.m_iteratorTuple) + if(mIteratorTuple < other.mIteratorTuple) return false; - return m_index >= other.m_index; + return mIndex >= other.mIndex; } #endif #pragma endregion comparisonoperators private: - IteratorTuple m_iteratorTuple; - IdxType m_index; + IdxType mIndex; + IteratorTuple mIteratorTuple; + + template + struct seq { }; + + template + struct gen_seq : gen_seq { }; + + template + struct gen_seq<0, Is...> : seq { }; + + template + auto makeValueTuple(std::tuple& t, seq) + -> std::tuple::type...> + { + return std::forward_as_tuple(*std::get(t)...); + } + + template + auto makeValueTuple(std::tuple& t) + -> std::tuple::type...> + { + return makeValueTuple(t, gen_seq()); + } template - inline typename std::enable_if::type - for_each(std::tuple &, FuncT) // Unused arguments are given no names - { - } + inline typename std::enable_if::type forEach(std::tuple &, FuncT) // Unused arguments are given no names + { + } template - inline typename std::enable_if::type - for_each(std::tuple& t, FuncT f) - { - f(std::get(t)); - for_each(t, f); - } + inline typename std::enable_if::type forEach(std::tuple& t, FuncT f) + { + f(std::get(t)); + forEach(t, f); + } }; } // namespace iterator From d8ff86ccbe87d60b43a67f113ffaeb7fe9388328 Mon Sep 17 00:00:00 2001 From: Victor Date: Tue, 25 Jan 2022 04:45:35 +0100 Subject: [PATCH 6/9] Fix ZipIterator and use deviceTransform for example, still need more checking --- example/zipIterator/src/zipIterator-main.cpp | 116 ++++++++++-------- include/vikunja/mem/iterator/ZipIterator.hpp | 119 ++++++++++++++----- 2 files changed, 150 insertions(+), 85 deletions(-) diff --git a/example/zipIterator/src/zipIterator-main.cpp b/example/zipIterator/src/zipIterator-main.cpp index f9f493f..61170b6 100644 --- a/example/zipIterator/src/zipIterator-main.cpp +++ b/example/zipIterator/src/zipIterator-main.cpp @@ -7,7 +7,7 @@ * file, You can obtain one at http://mozilla.org/MPL/2.0/. */ -#include +#include #include #include @@ -27,12 +27,12 @@ inline typename std::enable_if::type forEach(std::tuple forEach(t, f); } -template -void printTuple(IteratorTuple tuple) +template +void printTuple(IteratorTuplePtr tuple) { std::cout << "tuple("; int index = 0; - int tupleSize = std::tuple_size{}; + int tupleSize = std::tuple_size{}; forEach(tuple, [&index, tupleSize](auto &x) { std::cout << x << (++index < tupleSize ? ", " : ""); }); std::cout << ")"; } @@ -123,9 +123,10 @@ int main() std::cout << "\nTesting zip iterator in host with tuple\n\n"; - using IteratorTuple = std::tuple; - IteratorTuple zipTuple = std::make_tuple(hostNative, hostNativeChar, hostNativeDouble); - vikunja::mem::iterator::ZipIterator zipIter(zipTuple); + using IteratorTuplePtr = std::tuple; + using IteratorTupleVal = std::tuple; + IteratorTuplePtr zipTuple = std::make_tuple(hostNative, hostNativeChar, hostNativeDouble); + vikunja::mem::iterator::ZipIterator zipIter(zipTuple); std::cout << "*zipIter: "; printTuple(*zipIter); @@ -144,29 +145,36 @@ int main() std::cout << "\n\n"; zipIter += 6; - std::cout << "zipIter += 6;\n" + std::cout << "*zipIter += 6;\n" << "*zipIter: "; printTuple(*zipIter); std::cout << "\n\n"; zipIter -= 2; - std::cout << "zipIter -= 2;\n" + std::cout << "*zipIter -= 2;\n" << "*zipIter: "; printTuple(*zipIter); std::cout << "\n\n"; - std::cout << "--zipIter: "; + std::cout << "*--zipIter: "; printTuple(*--zipIter); std::cout << "\n*zipIter: "; printTuple(*zipIter); std::cout << "\n\n"; - std::cout << "zipIter--: "; + std::cout << "*zipIter--: "; printTuple(*zipIter--); std::cout << "\n*zipIter: "; printTuple(*zipIter); std::cout << "\n\n"; + std::cout << "Double the number values of the tuple:\n" + << "zipIter = std::make_tuple(2 * std::get<0>(*zipIter), std::get<1>(*zipIter), 2 * std::get<2>(*zipIter));\n" + << "*zipIter: "; + zipIter = std::make_tuple(2 * std::get<0>(*zipIter), std::get<1>(*zipIter), 2 * std::get<2>(*zipIter)); + printTuple(*zipIter); + std::cout << "\n\n"; + std::cout << "*(zipIter + 2): "; printTuple(*(zipIter + 2)); std::cout << "\n\n"; @@ -179,8 +187,12 @@ int main() printTuple(zipIter[0]); std::cout << "\n"; - std::cout << "zipIter[3]: "; - printTuple(zipIter[3]); + std::cout << "zipIter[2]: "; + printTuple(zipIter[2]); + std::cout << "\n"; + + std::cout << "zipIter[4] (number values has been doubled): "; + printTuple(zipIter[4]); std::cout << "\n"; std::cout << "zipIter[6]: "; @@ -189,45 +201,45 @@ int main() std::cout << "zipIter[9]: "; printTuple(zipIter[9]); - std::cout << "\n\n"; - - - - std::cout << "-----\n\n" - << "Failed when building for accelerator testing\n\n"; - - // std::cout << "Testing accelerator: " << alpaka::getAccName() << " with size: " << n << "\n\n"; - - // IteratorTuple deviceZipTuple = std::make_tuple(deviceNative, deviceNativeChar, deviceNativeDouble); - // vikunja::mem::iterator::ZipIterator deviceZipIter(deviceZipTuple); - - // using IteratorTupleReduceResult = std::tuple; - - // // Use Lambda function for reduction - // auto sum = [] ALPAKA_FN_HOST_ACC(IteratorTuple const i, IteratorTuple const j) - // { - // IteratorTupleReduceResult tmp = std::make_tuple(*std::get<0>(i) + *std::get<0>(j), *std::get<1>(i), *std::get<2>(i) + *std::get<2>(j)); - // return tmp; - // }; - // auto doubleNum = [] ALPAKA_FN_HOST_ACC(IteratorTuple const i) - // { - // IteratorTupleReduceResult tmp = std::make_tuple(2 * *std::get<0>(i), *std::get<1>(i), 2 * *std::get<2>(i)); - // return tmp; - // }; - - // // TRANSFORM_REDUCE CALL: - // // Takes the arguments: accelerator device, host device, accelerator queue, size of data, pointer-like to memory, - // // transform lambda, reduce lambda. - // auto transformReduceResult = vikunja::reduce::deviceTransformReduce( - // devAcc, - // devHost, - // queueAcc, - // n, - // deviceZipIter, - // doubleNum, - // sum); - - + std::cout << "\n\n" + << "-----\n\n"; + + // Revert the number values for index 4 + zipIter = std::make_tuple(std::get<0>(*zipIter) / 2, std::get<1>(*zipIter), std::get<2>(*zipIter) / 2); + + IteratorTuplePtr deviceZipTuple = std::make_tuple(deviceNative, deviceNativeChar, deviceNativeDouble); + vikunja::mem::iterator::ZipIterator deviceZipIter(deviceZipTuple); + + auto doubleNum = [] ALPAKA_FN_HOST_ACC(IteratorTupleVal const i) + { + return std::make_tuple(2 * std::get<0>(i), std::get<1>(i), 2 * std::get<2>(i)); + }; + + vikunja::transform::deviceTransform( + devAcc, + queueAcc, + extent[Dim::value - 1u], + deviceZipIter, + deviceZipIter, + doubleNum); + + // Copy the data back to the host for validation. + alpaka::memcpy(queueAcc, hostMem, deviceMem, extent); + + TRed resultSum = std::accumulate(hostNative, hostNative + extent.prod(), 0); + TRed expectedResult = extent.prod() * (extent.prod() + 1); + + std::cout << "Testing accelerator: " << alpaka::getAccName() << " with size: " << extent.prod() << "\n"; + if(expectedResult == resultSum) + { + std::cout << "Transform was successful!\n\n"; + } + else + { + std::cout << "Transform was not successful!\n" + << "expected result: " << expectedResult << "\n" + << "actual result: " << resultSum << "\n\n"; + } return 0; } diff --git a/include/vikunja/mem/iterator/ZipIterator.hpp b/include/vikunja/mem/iterator/ZipIterator.hpp index 62271cc..99fe9d8 100644 --- a/include/vikunja/mem/iterator/ZipIterator.hpp +++ b/include/vikunja/mem/iterator/ZipIterator.hpp @@ -40,52 +40,64 @@ namespace vikunja { /** * @brief A zip iterator that takes multiple input sequences and yields a sequence of tuples - * @tparam IteratorTuple The type of the data + * @tparam IteratorTuplePtr The type of the data + * @tparam IteratorTupleVal The type of the data * @tparam IdxType The type of the index */ - template + template class ZipIterator { public: // Need all 5 of these types for iterator_traits - using reference = IteratorTuple&; - using value_type = IteratorTuple; - using pointer = IteratorTuple*; + using reference = IteratorTupleVal&; + using value_type = IteratorTupleVal; + using pointer = IteratorTupleVal*; using difference_type = IdxType; using iterator_category = std::random_access_iterator_tag; /** * @brief Constructor for the ZipIterator - * @param iteratorTuple The tuple to initialize the iterator with + * @param iteratorTuplePtr The tuple to initialize the iterator with * @param idx The index for the iterator, default 0 */ - ZipIterator(IteratorTuple iteratorTuple, const IdxType& idx = static_cast(0)) + ZipIterator(IteratorTuplePtr iteratorTuplePtr, const IdxType& idx = static_cast(0)) : mIndex(idx) - , mIteratorTuple(iteratorTuple) + , mIteratorTuplePtr(iteratorTuplePtr) + , mIteratorTupleVal(makeValueTuple(mIteratorTuplePtr)) { if (idx != 0) - forEach(mIteratorTuple, [idx](auto &x) { x += idx; }); + { + forEach(mIteratorTuplePtr, [idx](auto &x) { x += idx; }); + mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); + } } /** * @brief Dereference operator to receive the stored value */ - NODISCARD ALPAKA_FN_INLINE auto operator*() + NODISCARD ALPAKA_FN_INLINE IteratorTupleVal& operator*() { - return makeValueTuple(mIteratorTuple); + return mIteratorTupleVal; } /** * @brief Index operator to get stored value at some given offset from this iterator */ - NODISCARD ALPAKA_FN_INLINE auto operator[](const IdxType idx) + NODISCARD ALPAKA_FN_INLINE const IteratorTupleVal operator[](const IdxType idx) { - IteratorTuple tmp = mIteratorTuple; + IteratorTuplePtr tmp = mIteratorTuplePtr; IdxType indexDiff = idx - mIndex; forEach(tmp, [indexDiff](auto &x) { x += indexDiff; }); return makeValueTuple(tmp); } + NODISCARD ALPAKA_FN_INLINE IteratorTupleVal& operator=(IteratorTupleVal iteratorTupleVal) + { + updateIteratorTupleValue(iteratorTupleVal); + mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); + return mIteratorTupleVal; + } + #pragma region arithmeticoperators /** * @brief Prefix increment operator @@ -93,7 +105,8 @@ namespace vikunja NODISCARD ALPAKA_FN_INLINE ZipIterator& operator++() { ++mIndex; - forEach(mIteratorTuple, [](auto &x) { ++x; }); + forEach(mIteratorTuplePtr, [](auto &x) { ++x; }); + mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); return *this; } @@ -105,7 +118,8 @@ namespace vikunja { ZipIterator tmp = *this; ++mIndex; - forEach(mIteratorTuple, [](auto &x) { ++x; }); + forEach(mIteratorTuplePtr, [](auto &x) { ++x; }); + mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); return tmp; } @@ -115,7 +129,8 @@ namespace vikunja NODISCARD ALPAKA_FN_INLINE ZipIterator& operator--() { --mIndex; - forEach(mIteratorTuple, [](auto &x) { --x; }); + forEach(mIteratorTuplePtr, [](auto &x) { --x; }); + mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); return *this; } @@ -127,39 +142,63 @@ namespace vikunja { ZipIterator tmp = *this; --mIndex; - forEach(mIteratorTuple, [](auto &x) { --x; }); + forEach(mIteratorTuplePtr, [](auto &x) { --x; }); + mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); return tmp; } /** * @brief Add an index to this iterator */ - NODISCARD ALPAKA_FN_INLINE ZipIterator operator+(const IdxType idx) + NODISCARD ALPAKA_FN_INLINE ZipIterator operator+(const int idx) { - IteratorTuple tmp = mIteratorTuple; + IteratorTuplePtr tmp = mIteratorTuplePtr; IdxType indexDiff = mIndex; forEach(tmp, [indexDiff](auto &x) { x -= indexDiff; }); return ZipIterator(tmp, mIndex + idx); } + /** + * @brief Add an index to this iterator + */ + NODISCARD friend ALPAKA_FN_INLINE ZipIterator operator+(ZipIterator zipIter, const IdxType idx) + { + IteratorTuplePtr tmp = zipIter.mIteratorTuplePtr; + IdxType indexDiff = zipIter.mIndex; + zipIter.forEach(tmp, [indexDiff](auto &x) { x -= indexDiff; }); + return ZipIterator(tmp, zipIter.mIndex + idx); + } + /** * @brief Subtract an index from this iterator */ - NODISCARD ALPAKA_FN_INLINE ZipIterator operator-(const IdxType idx) + NODISCARD ALPAKA_FN_INLINE ZipIterator operator-(const int idx) { - IteratorTuple tmp = mIteratorTuple; + IteratorTuplePtr tmp = mIteratorTuplePtr; IdxType indexDiff = mIndex; forEach(tmp, [indexDiff](auto &x) { x -= indexDiff; }); return ZipIterator(tmp, mIndex - idx); } + /** + * @brief Subtract an index from this iterator + */ + NODISCARD friend ALPAKA_FN_INLINE ZipIterator operator-(ZipIterator zipIter, const IdxType idx) + { + IteratorTuplePtr tmp = zipIter.mIteratorTuplePtr; + IdxType indexDiff = zipIter.mIndex; + zipIter.forEach(tmp, [indexDiff](auto &x) { x -= indexDiff; }); + return ZipIterator(tmp, zipIter.mIndex - idx); + } + /** * @brief Add an index to this iterator */ NODISCARD ALPAKA_FN_INLINE ZipIterator& operator+=(const IdxType idx) { mIndex += idx; - forEach(mIteratorTuple, [idx](auto &x) { x += idx; }); + forEach(mIteratorTuplePtr, [idx](auto &x) { x += idx; }); + mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); return *this; } @@ -169,7 +208,8 @@ namespace vikunja NODISCARD ALPAKA_FN_INLINE ZipIterator& operator-=(const IdxType idx) { mIndex -= idx; - forEach(mIteratorTuple, [idx](auto &x) { x -= idx; }); + forEach(mIteratorTuplePtr, [idx](auto &x) { x -= idx; }); + mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); return *this; } @@ -193,7 +233,7 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE bool operator==(const ZipIterator& other) const noexcept { - return mIteratorTuple == other.mIteratorTuple && mIndex == other.mIndex; + return mIteratorTuplePtr == other.mIteratorTuplePtr && mIndex == other.mIndex; } /** @@ -209,9 +249,9 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE bool operator<(const ZipIterator& other) const noexcept { - if(mIteratorTuple < other.mIteratorTuple) + if(mIteratorTuplePtr < other.mIteratorTuplePtr) return true; - if(mIteratorTuple > other.mIteratorTuple) + if(mIteratorTuplePtr > other.mIteratorTuplePtr) return false; return mIndex < other.mIndex; } @@ -221,9 +261,9 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE bool operator>(const ZipIterator& other) const noexcept { - if(mIteratorTuple > other.mIteratorTuple) + if(mIteratorTuplePtr > other.mIteratorTuplePtr) return true; - if(mIteratorTuple < other.mIteratorTuple) + if(mIteratorTuplePtr < other.mIteratorTuplePtr) return false; return mIndex > other.mIndex; } @@ -233,9 +273,9 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE bool operator<=(const ZipIterator& other) const noexcept { - if(mIteratorTuple < other.mIteratorTuple) + if(mIteratorTuplePtr < other.mIteratorTuplePtr) return true; - if(mIteratorTuple > other.mIteratorTuple) + if(mIteratorTuplePtr > other.mIteratorTuplePtr) return false; return mIndex <= other.mIndex; } @@ -245,9 +285,9 @@ namespace vikunja */ NODISCARD ALPAKA_FN_INLINE bool operator>=(const ZipIterator& other) const noexcept { - if(mIteratorTuple > other.mIteratorTuple) + if(mIteratorTuplePtr > other.mIteratorTuplePtr) return true; - if(mIteratorTuple < other.mIteratorTuple) + if(mIteratorTuplePtr < other.mIteratorTuplePtr) return false; return mIndex >= other.mIndex; } @@ -257,7 +297,8 @@ namespace vikunja private: IdxType mIndex; - IteratorTuple mIteratorTuple; + IteratorTuplePtr mIteratorTuplePtr; + IteratorTupleVal mIteratorTupleVal; template struct seq { }; @@ -293,6 +334,18 @@ namespace vikunja f(std::get(t)); forEach(t, f); } + + template + inline typename std::enable_if::type updateIteratorTupleValue(std::tuple &) // Unused arguments are given no names + { + } + + template + inline typename std::enable_if::type updateIteratorTupleValue(std::tuple& t) + { + *std::get(mIteratorTuplePtr) = std::get(t); + updateIteratorTupleValue(t); + } }; } // namespace iterator From f6625f8942292a6eaeb703dc89ec2f4f4aba1a16 Mon Sep 17 00:00:00 2001 From: Victor Date: Sun, 30 Jan 2022 21:05:37 +0100 Subject: [PATCH 7/9] Update ZipIterator example: - Change transform output from fancy iterator to a memory object of std::tuple (like the return type of the functor) - Change transform result check by comparing element by element - Directly use the concrete types Update ZipIterator.hpp: - Remove = operator --- example/zipIterator/src/zipIterator-main.cpp | 139 +++++++++++-------- include/vikunja/mem/iterator/ZipIterator.hpp | 125 ++++++----------- 2 files changed, 120 insertions(+), 144 deletions(-) diff --git a/example/zipIterator/src/zipIterator-main.cpp b/example/zipIterator/src/zipIterator-main.cpp index 61170b6..d93f850 100644 --- a/example/zipIterator/src/zipIterator-main.cpp +++ b/example/zipIterator/src/zipIterator-main.cpp @@ -27,12 +27,12 @@ inline typename std::enable_if::type forEach(std::tuple forEach(t, f); } -template -void printTuple(IteratorTuplePtr tuple) +template +void printTuple(IteratorTupleVal tuple) { - std::cout << "tuple("; + std::cout << "("; int index = 0; - int tupleSize = std::tuple_size{}; + int tupleSize = std::tuple_size{}; forEach(tuple, [&index, tupleSize](auto &x) { std::cout << x << (++index < tupleSize ? ", " : ""); }); std::cout << ")"; } @@ -42,11 +42,6 @@ int main() // Define the accelerator here. Must be one of the enabled accelerators. using TAcc = alpaka::AccCpuSerial, std::uint64_t>; - // Types of the data that will be reduced - using TRed = uint64_t; - using TRedChar = char; - using TRedDouble = double; - // Alpaka index type using Idx = alpaka::Idx; // Alpaka dimension type @@ -81,57 +76,52 @@ int main() QueueAcc queueAcc(devAcc); // allocate memory both on host and device. - auto deviceMem(alpaka::allocBuf(devAcc, extent)); - auto hostMem(alpaka::allocBuf(devHost, extent)); + auto deviceMem(alpaka::allocBuf(devAcc, extent)); + auto hostMem(alpaka::allocBuf(devHost, extent)); // Fill memory on host with numbers from 0...n-1. - TRed* hostNative = alpaka::getPtrNative(hostMem); + uint64_t* hostNative = alpaka::getPtrNative(hostMem); for(Idx i = 0; i < n; ++i) - hostNative[i] = static_cast(i + 1); + hostNative[i] = static_cast(i + 1); // Copy to accelerator. alpaka::memcpy(queueAcc, deviceMem, hostMem, extent); - TRed* deviceNative = alpaka::getPtrNative(deviceMem); + uint64_t* deviceNative = alpaka::getPtrNative(deviceMem); // allocate memory both on host and device. - auto deviceMemChar(alpaka::allocBuf(devAcc, extent)); - auto hostMemChar(alpaka::allocBuf(devHost, extent)); + auto deviceMemChar(alpaka::allocBuf(devAcc, extent)); + auto hostMemChar(alpaka::allocBuf(devHost, extent)); + std::vector chars = { 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j' }; // Fill memory on host with char from 'a' to 'j'. - TRedChar* hostNativeChar = alpaka::getPtrNative(hostMemChar); - hostNativeChar[0] = 'a'; - hostNativeChar[1] = 'b'; - hostNativeChar[2] = 'c'; - hostNativeChar[3] = 'd'; - hostNativeChar[4] = 'e'; - hostNativeChar[5] = 'f'; - hostNativeChar[6] = 'g'; - hostNativeChar[7] = 'h'; - hostNativeChar[8] = 'i'; - hostNativeChar[9] = 'j'; + char* hostNativeChar = alpaka::getPtrNative(hostMemChar); + for(Idx i = 0; i < n; ++i) + { + hostNativeChar[i] = chars[i]; + } // Copy to accelerator. alpaka::memcpy(queueAcc, deviceMemChar, hostMemChar, extent); - TRedChar* deviceNativeChar = alpaka::getPtrNative(deviceMemChar); + char* deviceNativeChar = alpaka::getPtrNative(deviceMemChar); // allocate memory both on host and device. - auto deviceMemDouble(alpaka::allocBuf(devAcc, extent)); - auto hostMemDouble(alpaka::allocBuf(devHost, extent)); + auto deviceMemDouble(alpaka::allocBuf(devAcc, extent)); + auto hostMemDouble(alpaka::allocBuf(devHost, extent)); // Fill memory on host with double numbers from 10.12...(n-1 + 10.12). - TRedDouble* hostNativeDouble = alpaka::getPtrNative(hostMemDouble); + double* hostNativeDouble = alpaka::getPtrNative(hostMemDouble); for(Idx i = 0; i < n; ++i) - hostNativeDouble[i] = static_cast(i + 10.12); + hostNativeDouble[i] = static_cast(i + 10.12); // Copy to accelerator. alpaka::memcpy(queueAcc, deviceMemDouble, hostMemDouble, extent); - TRedDouble* deviceNativeDouble = alpaka::getPtrNative(deviceMemDouble); + double* deviceNativeDouble = alpaka::getPtrNative(deviceMemDouble); std::cout << "\nTesting zip iterator in host with tuple\n\n"; - using IteratorTuplePtr = std::tuple; - using IteratorTupleVal = std::tuple; + using IteratorTuplePtr = std::tuple; + using IteratorTupleVal = std::tuple; IteratorTuplePtr zipTuple = std::make_tuple(hostNative, hostNativeChar, hostNativeDouble); vikunja::mem::iterator::ZipIterator zipIter(zipTuple); std::cout << "*zipIter: "; printTuple(*zipIter); std::cout << "\n\n"; - + std::cout << "*++zipIter: "; printTuple(*++zipIter); std::cout << "\n*zipIter: "; @@ -168,12 +158,12 @@ int main() printTuple(*zipIter); std::cout << "\n\n"; - std::cout << "Double the number values of the tuple:\n" - << "zipIter = std::make_tuple(2 * std::get<0>(*zipIter), std::get<1>(*zipIter), 2 * std::get<2>(*zipIter));\n" - << "*zipIter: "; - zipIter = std::make_tuple(2 * std::get<0>(*zipIter), std::get<1>(*zipIter), 2 * std::get<2>(*zipIter)); - printTuple(*zipIter); - std::cout << "\n\n"; + // std::cout << "Double the number values of the tuple:\n" + // << "zipIter = std::make_tuple(2 * std::get<0>(*zipIter), std::get<1>(*zipIter), 2 * std::get<2>(*zipIter));\n" + // << "*zipIter: "; + // zipIter = std::make_tuple(2 * std::get<0>(*zipIter), std::get<1>(*zipIter), 2 * std::get<2>(*zipIter)); + // printTuple(*zipIter); + // std::cout << "\n\n"; std::cout << "*(zipIter + 2): "; printTuple(*(zipIter + 2)); @@ -191,28 +181,39 @@ int main() printTuple(zipIter[2]); std::cout << "\n"; - std::cout << "zipIter[4] (number values has been doubled): "; + // std::cout << "zipIter[4] (number values has been doubled): "; + // printTuple(zipIter[4]); + // std::cout << "\n"; + + // std::cout << "Revert the number values for index 4\n"; + // zipIter = std::make_tuple(std::get<0>(*zipIter) / 2, std::get<1>(*zipIter), std::get<2>(*zipIter) / 2); + + std::cout << "zipIter[4]: "; printTuple(zipIter[4]); std::cout << "\n"; std::cout << "zipIter[6]: "; printTuple(zipIter[6]); std::cout << "\n"; - + std::cout << "zipIter[9]: "; printTuple(zipIter[9]); std::cout << "\n\n" << "-----\n\n"; - // Revert the number values for index 4 - zipIter = std::make_tuple(std::get<0>(*zipIter) / 2, std::get<1>(*zipIter), std::get<2>(*zipIter) / 2); - IteratorTuplePtr deviceZipTuple = std::make_tuple(deviceNative, deviceNativeChar, deviceNativeDouble); vikunja::mem::iterator::ZipIterator deviceZipIter(deviceZipTuple); - auto doubleNum = [] ALPAKA_FN_HOST_ACC(IteratorTupleVal const i) + auto deviceMemResult(alpaka::allocBuf(devAcc, extent)); + auto hostMemResult(alpaka::allocBuf(devHost, extent)); + IteratorTupleVal* hostNativeResultPtr = alpaka::getPtrNative(hostMemResult); + IteratorTupleVal* deviceNativeResultPtr = alpaka::getPtrNative(deviceMemResult); + + auto doubleNum = [] ALPAKA_FN_HOST_ACC(IteratorTupleVal const& t) { - return std::make_tuple(2 * std::get<0>(i), std::get<1>(i), 2 * std::get<2>(i)); + // return std::make_tuple(2 * std::get<0>(t), std::get<1>(t), 2 * std::get<2>(t)); + // return std::make_tuple(static_cast(5), 'e', static_cast(14.12)); + return t; }; vikunja::transform::deviceTransform( @@ -220,26 +221,42 @@ int main() queueAcc, extent[Dim::value - 1u], deviceZipIter, - deviceZipIter, + deviceNativeResultPtr, doubleNum); // Copy the data back to the host for validation. - alpaka::memcpy(queueAcc, hostMem, deviceMem, extent); + alpaka::memcpy(queueAcc, hostMemResult, deviceMemResult, extent); - TRed resultSum = std::accumulate(hostNative, hostNative + extent.prod(), 0); - TRed expectedResult = extent.prod() * (extent.prod() + 1); + std::cout << "Testing accelerator: " << alpaka::getAccName() << " with size: " << extent.prod() << "\n" + << "-----\n"; - std::cout << "Testing accelerator: " << alpaka::getAccName() << " with size: " << extent.prod() << "\n"; - if(expectedResult == resultSum) + bool isTransformSuccess = true; + for(Idx i = 0; i < n; ++i) { - std::cout << "Transform was successful!\n\n"; + std::cout << "n=" << i << " | Expected result: (" + << 2 * (i + 1) << ", " << chars[i] << ", " << 2 * (i + 10.12) + << ") | Actual result: "; + printTuple(hostNativeResultPtr[i]); + + if((2 * (i + 1) == std::get<0>(hostNativeResultPtr[i])) && + (chars[i] == std::get<1>(hostNativeResultPtr[i])) && + (2 * (i + 10.12) == std::get<2>(hostNativeResultPtr[i])) + ) { + std::cout << " | OK\n"; + } + else + { + std::cout << " | NOT OK\n"; + isTransformSuccess = false; + } } + + if(isTransformSuccess) + std::cout << "-----\n" + << "Transform was successful!\n\n"; else - { - std::cout << "Transform was not successful!\n" - << "expected result: " << expectedResult << "\n" - << "actual result: " << resultSum << "\n\n"; - } + std::cout << "-----\n" + << "Transform was NOT successful!\n\n"; return 0; } diff --git a/include/vikunja/mem/iterator/ZipIterator.hpp b/include/vikunja/mem/iterator/ZipIterator.hpp index 99fe9d8..ac1e6d0 100644 --- a/include/vikunja/mem/iterator/ZipIterator.hpp +++ b/include/vikunja/mem/iterator/ZipIterator.hpp @@ -60,7 +60,7 @@ namespace vikunja * @param iteratorTuplePtr The tuple to initialize the iterator with * @param idx The index for the iterator, default 0 */ - ZipIterator(IteratorTuplePtr iteratorTuplePtr, const IdxType& idx = static_cast(0)) + constexpr ZipIterator(IteratorTuplePtr iteratorTuplePtr, const IdxType& idx = static_cast(0)) : mIndex(idx) , mIteratorTuplePtr(iteratorTuplePtr) , mIteratorTupleVal(makeValueTuple(mIteratorTuplePtr)) @@ -75,7 +75,7 @@ namespace vikunja /** * @brief Dereference operator to receive the stored value */ - NODISCARD ALPAKA_FN_INLINE IteratorTupleVal& operator*() + NODISCARD constexpr ALPAKA_FN_INLINE IteratorTupleVal& operator*() { return mIteratorTupleVal; } @@ -83,7 +83,7 @@ namespace vikunja /** * @brief Index operator to get stored value at some given offset from this iterator */ - NODISCARD ALPAKA_FN_INLINE const IteratorTupleVal operator[](const IdxType idx) + NODISCARD constexpr ALPAKA_FN_INLINE const IteratorTupleVal operator[](const IdxType idx) { IteratorTuplePtr tmp = mIteratorTuplePtr; IdxType indexDiff = idx - mIndex; @@ -91,18 +91,18 @@ namespace vikunja return makeValueTuple(tmp); } - NODISCARD ALPAKA_FN_INLINE IteratorTupleVal& operator=(IteratorTupleVal iteratorTupleVal) - { - updateIteratorTupleValue(iteratorTupleVal); - mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); - return mIteratorTupleVal; - } + // NODISCARD constexpr ALPAKA_FN_INLINE IteratorTupleVal& operator=(IteratorTupleVal iteratorTupleVal) + // { + // updateIteratorTupleValue(iteratorTupleVal); + // mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); + // return mIteratorTupleVal; + // } #pragma region arithmeticoperators /** * @brief Prefix increment operator */ - NODISCARD ALPAKA_FN_INLINE ZipIterator& operator++() + constexpr ALPAKA_FN_INLINE ZipIterator& operator++() { ++mIndex; forEach(mIteratorTuplePtr, [](auto &x) { ++x; }); @@ -114,7 +114,7 @@ namespace vikunja * @brief Postfix increment operator * @note Use prefix increment operator instead if possible to avoid copies */ - ALPAKA_FN_INLINE ZipIterator operator++(int) + constexpr ZipIterator operator++(int) { ZipIterator tmp = *this; ++mIndex; @@ -126,7 +126,7 @@ namespace vikunja /** * @brief Prefix decrement operator */ - NODISCARD ALPAKA_FN_INLINE ZipIterator& operator--() + constexpr ALPAKA_FN_INLINE ZipIterator& operator--() { --mIndex; forEach(mIteratorTuplePtr, [](auto &x) { --x; }); @@ -138,7 +138,7 @@ namespace vikunja * @brief Postfix decrement operator * @note Use prefix decrement operator instead if possible to avoid copies */ - ALPAKA_FN_INLINE ZipIterator operator--(int) + constexpr ALPAKA_FN_INLINE ZipIterator operator--(int) { ZipIterator tmp = *this; --mIndex; @@ -150,51 +150,25 @@ namespace vikunja /** * @brief Add an index to this iterator */ - NODISCARD ALPAKA_FN_INLINE ZipIterator operator+(const int idx) - { - IteratorTuplePtr tmp = mIteratorTuplePtr; - IdxType indexDiff = mIndex; - forEach(tmp, [indexDiff](auto &x) { x -= indexDiff; }); - return ZipIterator(tmp, mIndex + idx); - } - - /** - * @brief Add an index to this iterator - */ - NODISCARD friend ALPAKA_FN_INLINE ZipIterator operator+(ZipIterator zipIter, const IdxType idx) - { - IteratorTuplePtr tmp = zipIter.mIteratorTuplePtr; - IdxType indexDiff = zipIter.mIndex; - zipIter.forEach(tmp, [indexDiff](auto &x) { x -= indexDiff; }); - return ZipIterator(tmp, zipIter.mIndex + idx); - } - - /** - * @brief Subtract an index from this iterator - */ - NODISCARD ALPAKA_FN_INLINE ZipIterator operator-(const int idx) + NODISCARD constexpr friend ALPAKA_FN_INLINE ZipIterator operator+(ZipIterator zipIter, const IdxType idx) { - IteratorTuplePtr tmp = mIteratorTuplePtr; - IdxType indexDiff = mIndex; - forEach(tmp, [indexDiff](auto &x) { x -= indexDiff; }); - return ZipIterator(tmp, mIndex - idx); + zipIter += idx; + return zipIter; } /** * @brief Subtract an index from this iterator */ - NODISCARD friend ALPAKA_FN_INLINE ZipIterator operator-(ZipIterator zipIter, const IdxType idx) + NODISCARD constexpr friend ALPAKA_FN_INLINE ZipIterator operator-(ZipIterator zipIter, const IdxType idx) { - IteratorTuplePtr tmp = zipIter.mIteratorTuplePtr; - IdxType indexDiff = zipIter.mIndex; - zipIter.forEach(tmp, [indexDiff](auto &x) { x -= indexDiff; }); - return ZipIterator(tmp, zipIter.mIndex - idx); + zipIter -= idx; + return zipIter; } /** * @brief Add an index to this iterator */ - NODISCARD ALPAKA_FN_INLINE ZipIterator& operator+=(const IdxType idx) + constexpr ALPAKA_FN_INLINE ZipIterator& operator+=(const IdxType idx) { mIndex += idx; forEach(mIteratorTuplePtr, [idx](auto &x) { x += idx; }); @@ -205,7 +179,7 @@ namespace vikunja /** * @brief Subtract an index from this iterator */ - NODISCARD ALPAKA_FN_INLINE ZipIterator& operator-=(const IdxType idx) + constexpr ALPAKA_FN_INLINE ZipIterator& operator-=(const IdxType idx) { mIndex -= idx; forEach(mIteratorTuplePtr, [idx](auto &x) { x -= idx; }); @@ -217,79 +191,64 @@ namespace vikunja #pragma region comparisonoperators -// if spaceship operator is available is being used we can use spaceship operator magic #ifdef USESPACESHIP /** * @brief Spaceship operator for comparisons */ - NODISCARD ALPAKA_FN_INLINE auto operator<=>(const ZipIterator& other) const noexcept = default; + NODISCARD constexpr ALPAKA_FN_INLINE auto operator<=>(const ZipIterator& other) const noexcept + { + return mIteratorTuplePtr.operator<=>(other.mIteratorTuplePtr); + } -// if cpp20 *isn't* defined we get to write 70 lines of boilerplate #else /** - * @brief Equality comparison, returns true if the iterators are the same + * @brief Equality comparison, returns true if the index are the same */ - NODISCARD ALPAKA_FN_INLINE bool operator==(const ZipIterator& other) const noexcept + NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator==(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - return mIteratorTuplePtr == other.mIteratorTuplePtr && mIndex == other.mIndex; + return zipIter.mIndex == other.mIndex; } /** * @brief Inequality comparison, negated equality operator */ - NODISCARD ALPAKA_FN_INLINE bool operator!=(const ZipIterator& other) const noexcept + NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator!=(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - return !operator==(other); + return !operator==(zipIter, other); } /** - * @brief Less than comparison, value is checked first, then index + * @brief Less than comparison, index is checked */ - NODISCARD ALPAKA_FN_INLINE bool operator<(const ZipIterator& other) const noexcept + NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator<(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - if(mIteratorTuplePtr < other.mIteratorTuplePtr) - return true; - if(mIteratorTuplePtr > other.mIteratorTuplePtr) - return false; - return mIndex < other.mIndex; + return zipIter.mIndex < other.mIndex; } /** - * @brief Greater than comparison, value is checked first, then index + * @brief Greater than comparison, index is checked */ - NODISCARD ALPAKA_FN_INLINE bool operator>(const ZipIterator& other) const noexcept + NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator>(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - if(mIteratorTuplePtr > other.mIteratorTuplePtr) - return true; - if(mIteratorTuplePtr < other.mIteratorTuplePtr) - return false; - return mIndex > other.mIndex; + return zipIter.mIndex > other.mIndex; } /** - * @brief Less than or equal comparison, value is checked first, then index + * @brief Less than or equal comparison, index is checked */ - NODISCARD ALPAKA_FN_INLINE bool operator<=(const ZipIterator& other) const noexcept + NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator<=(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - if(mIteratorTuplePtr < other.mIteratorTuplePtr) - return true; - if(mIteratorTuplePtr > other.mIteratorTuplePtr) - return false; - return mIndex <= other.mIndex; + return zipIter.mIndex <= other.mIndex; } /** - * @brief Greater than or equal comparison, value is checked first, then index + * @brief Greater than or equal comparison, index is checked */ - NODISCARD ALPAKA_FN_INLINE bool operator>=(const ZipIterator& other) const noexcept + NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator>=(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - if(mIteratorTuplePtr > other.mIteratorTuplePtr) - return true; - if(mIteratorTuplePtr < other.mIteratorTuplePtr) - return false; - return mIndex >= other.mIndex; + return zipIter.mIndex >= other.mIndex; } #endif From 53c7f1019d86f491cad6f8761840125f08c81146 Mon Sep 17 00:00:00 2001 From: Victor Date: Mon, 31 Jan 2022 16:11:56 +0100 Subject: [PATCH 8/9] Update ZipIterator: - Fix operator[] - Change constexpr to ALPAKA_FN_HOST_ACC --- example/zipIterator/src/zipIterator-main.cpp | 32 ++-- include/vikunja/mem/iterator/ZipIterator.hpp | 150 +++++++++---------- 2 files changed, 89 insertions(+), 93 deletions(-) diff --git a/example/zipIterator/src/zipIterator-main.cpp b/example/zipIterator/src/zipIterator-main.cpp index d93f850..c7e6caa 100644 --- a/example/zipIterator/src/zipIterator-main.cpp +++ b/example/zipIterator/src/zipIterator-main.cpp @@ -27,12 +27,12 @@ inline typename std::enable_if::type forEach(std::tuple forEach(t, f); } -template -void printTuple(IteratorTupleVal tuple) +template +void printTuple(TIteratorTupleVal tuple) { std::cout << "("; int index = 0; - int tupleSize = std::tuple_size{}; + int tupleSize = std::tuple_size{}; forEach(tuple, [&index, tupleSize](auto &x) { std::cout << x << (++index < tupleSize ? ", " : ""); }); std::cout << ")"; } @@ -113,10 +113,10 @@ int main() std::cout << "\nTesting zip iterator in host with tuple\n\n"; - using IteratorTuplePtr = std::tuple; - using IteratorTupleVal = std::tuple; - IteratorTuplePtr zipTuple = std::make_tuple(hostNative, hostNativeChar, hostNativeDouble); - vikunja::mem::iterator::ZipIterator zipIter(zipTuple); + using TIteratorTuplePtr = std::tuple; + using TIteratorTupleVal = std::tuple; + TIteratorTuplePtr zipTuple = std::make_tuple(hostNative, hostNativeChar, hostNativeDouble); + vikunja::mem::iterator::ZipIterator zipIter(zipTuple); std::cout << "*zipIter: "; printTuple(*zipIter); @@ -201,19 +201,17 @@ int main() std::cout << "\n\n" << "-----\n\n"; - IteratorTuplePtr deviceZipTuple = std::make_tuple(deviceNative, deviceNativeChar, deviceNativeDouble); - vikunja::mem::iterator::ZipIterator deviceZipIter(deviceZipTuple); + TIteratorTuplePtr deviceZipTuple = std::make_tuple(deviceNative, deviceNativeChar, deviceNativeDouble); + vikunja::mem::iterator::ZipIterator deviceZipIter(deviceZipTuple); - auto deviceMemResult(alpaka::allocBuf(devAcc, extent)); - auto hostMemResult(alpaka::allocBuf(devHost, extent)); - IteratorTupleVal* hostNativeResultPtr = alpaka::getPtrNative(hostMemResult); - IteratorTupleVal* deviceNativeResultPtr = alpaka::getPtrNative(deviceMemResult); + auto deviceMemResult(alpaka::allocBuf(devAcc, extent)); + auto hostMemResult(alpaka::allocBuf(devHost, extent)); + TIteratorTupleVal* hostNativeResultPtr = alpaka::getPtrNative(hostMemResult); + TIteratorTupleVal* deviceNativeResultPtr = alpaka::getPtrNative(deviceMemResult); - auto doubleNum = [] ALPAKA_FN_HOST_ACC(IteratorTupleVal const& t) + auto doubleNum = [] ALPAKA_FN_HOST_ACC(TIteratorTupleVal const& t) { - // return std::make_tuple(2 * std::get<0>(t), std::get<1>(t), 2 * std::get<2>(t)); - // return std::make_tuple(static_cast(5), 'e', static_cast(14.12)); - return t; + return std::make_tuple(2 * std::get<0>(t), std::get<1>(t), 2 * std::get<2>(t)); }; vikunja::transform::deviceTransform( diff --git a/include/vikunja/mem/iterator/ZipIterator.hpp b/include/vikunja/mem/iterator/ZipIterator.hpp index ac1e6d0..e918ad9 100644 --- a/include/vikunja/mem/iterator/ZipIterator.hpp +++ b/include/vikunja/mem/iterator/ZipIterator.hpp @@ -40,19 +40,19 @@ namespace vikunja { /** * @brief A zip iterator that takes multiple input sequences and yields a sequence of tuples - * @tparam IteratorTuplePtr The type of the data - * @tparam IteratorTupleVal The type of the data - * @tparam IdxType The type of the index + * @tparam TIteratorTuplePtr The type of the data + * @tparam TIteratorTupleVal The type of the data + * @tparam TIdx The type of the index */ - template + template class ZipIterator { public: // Need all 5 of these types for iterator_traits - using reference = IteratorTupleVal&; - using value_type = IteratorTupleVal; - using pointer = IteratorTupleVal*; - using difference_type = IdxType; + using reference = TIteratorTupleVal&; + using value_type = TIteratorTupleVal; + using pointer = TIteratorTupleVal*; + using difference_type = TIdx; using iterator_category = std::random_access_iterator_tag; /** @@ -60,53 +60,51 @@ namespace vikunja * @param iteratorTuplePtr The tuple to initialize the iterator with * @param idx The index for the iterator, default 0 */ - constexpr ZipIterator(IteratorTuplePtr iteratorTuplePtr, const IdxType& idx = static_cast(0)) - : mIndex(idx) - , mIteratorTuplePtr(iteratorTuplePtr) - , mIteratorTupleVal(makeValueTuple(mIteratorTuplePtr)) + ALPAKA_FN_HOST_ACC ZipIterator(TIteratorTuplePtr iteratorTuplePtr, const TIdx& idx = static_cast(0)) + : m_index(idx) + , m_iteratorTuplePtr(iteratorTuplePtr) + , m_iteratorTupleVal(makeValueTuple(m_iteratorTuplePtr)) { if (idx != 0) { - forEach(mIteratorTuplePtr, [idx](auto &x) { x += idx; }); - mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); + forEach(m_iteratorTuplePtr, [idx](auto &x) { x += idx; }); + m_iteratorTupleVal = makeValueTuple(m_iteratorTuplePtr); } } /** * @brief Dereference operator to receive the stored value */ - NODISCARD constexpr ALPAKA_FN_INLINE IteratorTupleVal& operator*() + NODISCARD ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE TIteratorTupleVal& operator*() { - return mIteratorTupleVal; + return m_iteratorTupleVal; } /** * @brief Index operator to get stored value at some given offset from this iterator */ - NODISCARD constexpr ALPAKA_FN_INLINE const IteratorTupleVal operator[](const IdxType idx) + NODISCARD ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE TIteratorTupleVal operator[](const TIdx idx) const { - IteratorTuplePtr tmp = mIteratorTuplePtr; - IdxType indexDiff = idx - mIndex; - forEach(tmp, [indexDiff](auto &x) { x += indexDiff; }); - return makeValueTuple(tmp); + TIdx indexDiff = idx - m_index; + return (*this + indexDiff).operator*(); } - // NODISCARD constexpr ALPAKA_FN_INLINE IteratorTupleVal& operator=(IteratorTupleVal iteratorTupleVal) + // NODISCARD ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE TIteratorTupleVal& operator=(TIteratorTupleVal iteratorTupleVal) // { // updateIteratorTupleValue(iteratorTupleVal); - // mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); - // return mIteratorTupleVal; + // m_iteratorTupleVal = makeValueTuple(m_iteratorTuplePtr); + // return m_iteratorTupleVal; // } #pragma region arithmeticoperators /** * @brief Prefix increment operator */ - constexpr ALPAKA_FN_INLINE ZipIterator& operator++() + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE ZipIterator& operator++() { - ++mIndex; - forEach(mIteratorTuplePtr, [](auto &x) { ++x; }); - mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); + ++m_index; + forEach(m_iteratorTuplePtr, [](auto &x) { ++x; }); + m_iteratorTupleVal = makeValueTuple(m_iteratorTuplePtr); return *this; } @@ -114,23 +112,23 @@ namespace vikunja * @brief Postfix increment operator * @note Use prefix increment operator instead if possible to avoid copies */ - constexpr ZipIterator operator++(int) + ALPAKA_FN_HOST_ACC ZipIterator operator++(int) { ZipIterator tmp = *this; - ++mIndex; - forEach(mIteratorTuplePtr, [](auto &x) { ++x; }); - mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); + ++m_index; + forEach(m_iteratorTuplePtr, [](auto &x) { ++x; }); + m_iteratorTupleVal = makeValueTuple(m_iteratorTuplePtr); return tmp; } /** * @brief Prefix decrement operator */ - constexpr ALPAKA_FN_INLINE ZipIterator& operator--() + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE ZipIterator& operator--() { - --mIndex; - forEach(mIteratorTuplePtr, [](auto &x) { --x; }); - mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); + --m_index; + forEach(m_iteratorTuplePtr, [](auto &x) { --x; }); + m_iteratorTupleVal = makeValueTuple(m_iteratorTuplePtr); return *this; } @@ -138,19 +136,19 @@ namespace vikunja * @brief Postfix decrement operator * @note Use prefix decrement operator instead if possible to avoid copies */ - constexpr ALPAKA_FN_INLINE ZipIterator operator--(int) + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE ZipIterator operator--(int) { ZipIterator tmp = *this; - --mIndex; - forEach(mIteratorTuplePtr, [](auto &x) { --x; }); - mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); + --m_index; + forEach(m_iteratorTuplePtr, [](auto &x) { --x; }); + m_iteratorTupleVal = makeValueTuple(m_iteratorTuplePtr); return tmp; } /** * @brief Add an index to this iterator */ - NODISCARD constexpr friend ALPAKA_FN_INLINE ZipIterator operator+(ZipIterator zipIter, const IdxType idx) + NODISCARD ALPAKA_FN_HOST_ACC friend ALPAKA_FN_INLINE ZipIterator operator+(ZipIterator zipIter, const TIdx idx) { zipIter += idx; return zipIter; @@ -159,7 +157,7 @@ namespace vikunja /** * @brief Subtract an index from this iterator */ - NODISCARD constexpr friend ALPAKA_FN_INLINE ZipIterator operator-(ZipIterator zipIter, const IdxType idx) + NODISCARD ALPAKA_FN_HOST_ACC friend ALPAKA_FN_INLINE ZipIterator operator-(ZipIterator zipIter, const TIdx idx) { zipIter -= idx; return zipIter; @@ -168,22 +166,22 @@ namespace vikunja /** * @brief Add an index to this iterator */ - constexpr ALPAKA_FN_INLINE ZipIterator& operator+=(const IdxType idx) + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE ZipIterator& operator+=(const TIdx idx) { - mIndex += idx; - forEach(mIteratorTuplePtr, [idx](auto &x) { x += idx; }); - mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); + m_index += idx; + forEach(m_iteratorTuplePtr, [idx](auto &x) { x += idx; }); + m_iteratorTupleVal = makeValueTuple(m_iteratorTuplePtr); return *this; } /** * @brief Subtract an index from this iterator */ - constexpr ALPAKA_FN_INLINE ZipIterator& operator-=(const IdxType idx) + ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE ZipIterator& operator-=(const TIdx idx) { - mIndex -= idx; - forEach(mIteratorTuplePtr, [idx](auto &x) { x -= idx; }); - mIteratorTupleVal = makeValueTuple(mIteratorTuplePtr); + m_index -= idx; + forEach(m_iteratorTuplePtr, [idx](auto &x) { x -= idx; }); + m_iteratorTupleVal = makeValueTuple(m_iteratorTuplePtr); return *this; } @@ -196,9 +194,9 @@ namespace vikunja /** * @brief Spaceship operator for comparisons */ - NODISCARD constexpr ALPAKA_FN_INLINE auto operator<=>(const ZipIterator& other) const noexcept + NODISCARD ALPAKA_FN_HOST_ACC ALPAKA_FN_INLINE auto operator<=>(const ZipIterator& other) const noexcept { - return mIteratorTuplePtr.operator<=>(other.mIteratorTuplePtr); + return m_iteratorTuplePtr.operator<=>(other.m_iteratorTuplePtr); } #else @@ -206,15 +204,15 @@ namespace vikunja /** * @brief Equality comparison, returns true if the index are the same */ - NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator==(const ZipIterator& zipIter, const ZipIterator& other) noexcept + NODISCARD ALPAKA_FN_HOST_ACC friend ALPAKA_FN_INLINE bool operator==(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - return zipIter.mIndex == other.mIndex; + return zipIter.m_index == other.m_index; } /** * @brief Inequality comparison, negated equality operator */ - NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator!=(const ZipIterator& zipIter, const ZipIterator& other) noexcept + NODISCARD ALPAKA_FN_HOST_ACC friend ALPAKA_FN_INLINE bool operator!=(const ZipIterator& zipIter, const ZipIterator& other) noexcept { return !operator==(zipIter, other); } @@ -222,42 +220,42 @@ namespace vikunja /** * @brief Less than comparison, index is checked */ - NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator<(const ZipIterator& zipIter, const ZipIterator& other) noexcept + NODISCARD ALPAKA_FN_HOST_ACC friend ALPAKA_FN_INLINE bool operator<(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - return zipIter.mIndex < other.mIndex; + return zipIter.m_index < other.m_index; } /** * @brief Greater than comparison, index is checked */ - NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator>(const ZipIterator& zipIter, const ZipIterator& other) noexcept + NODISCARD ALPAKA_FN_HOST_ACC friend ALPAKA_FN_INLINE bool operator>(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - return zipIter.mIndex > other.mIndex; + return zipIter.m_index > other.m_index; } /** * @brief Less than or equal comparison, index is checked */ - NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator<=(const ZipIterator& zipIter, const ZipIterator& other) noexcept + NODISCARD ALPAKA_FN_HOST_ACC friend ALPAKA_FN_INLINE bool operator<=(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - return zipIter.mIndex <= other.mIndex; + return zipIter.m_index <= other.m_index; } /** * @brief Greater than or equal comparison, index is checked */ - NODISCARD constexpr friend ALPAKA_FN_INLINE bool operator>=(const ZipIterator& zipIter, const ZipIterator& other) noexcept + NODISCARD ALPAKA_FN_HOST_ACC friend ALPAKA_FN_INLINE bool operator>=(const ZipIterator& zipIter, const ZipIterator& other) noexcept { - return zipIter.mIndex >= other.mIndex; + return zipIter.m_index >= other.m_index; } #endif #pragma endregion comparisonoperators private: - IdxType mIndex; - IteratorTuplePtr mIteratorTuplePtr; - IteratorTupleVal mIteratorTupleVal; + TIdx m_index; + TIteratorTuplePtr m_iteratorTuplePtr; + TIteratorTupleVal m_iteratorTupleVal; template struct seq { }; @@ -294,17 +292,17 @@ namespace vikunja forEach(t, f); } - template - inline typename std::enable_if::type updateIteratorTupleValue(std::tuple &) // Unused arguments are given no names - { - } + // template + // inline typename std::enable_if::type updateIteratorTupleValue(std::tuple &) // Unused arguments are given no names + // { + // } - template - inline typename std::enable_if::type updateIteratorTupleValue(std::tuple& t) - { - *std::get(mIteratorTuplePtr) = std::get(t); - updateIteratorTupleValue(t); - } + // template + // inline typename std::enable_if::type updateIteratorTupleValue(std::tuple& t) + // { + // *std::get(m_iteratorTuplePtr) = std::get(t); + // updateIteratorTupleValue(t); + // } }; } // namespace iterator From 06ef6671ee7488bd42f3bdfae14899c4d03d2028 Mon Sep 17 00:00:00 2001 From: Victor Date: Mon, 31 Jan 2022 18:01:42 +0100 Subject: [PATCH 9/9] Add ZipIterator benchmark --- include/vikunja/mem/iterator/ZipIterator.hpp | 5 + test/benchmarks/transform/CMakeLists.txt | 18 +++ .../bench_vikunja_transform_zipiter.cpp | 108 ++++++++++++++++++ 3 files changed, 131 insertions(+) create mode 100644 test/benchmarks/transform/bench_vikunja_transform_zipiter.cpp diff --git a/include/vikunja/mem/iterator/ZipIterator.hpp b/include/vikunja/mem/iterator/ZipIterator.hpp index e918ad9..1999769 100644 --- a/include/vikunja/mem/iterator/ZipIterator.hpp +++ b/include/vikunja/mem/iterator/ZipIterator.hpp @@ -163,6 +163,11 @@ namespace vikunja return zipIter; } + NODISCARD ALPAKA_FN_HOST_ACC friend ALPAKA_FN_INLINE TIdx operator-(const ZipIterator& zipIter, const ZipIterator& other) + { + return zipIter.m_index - other.m_index; + } + /** * @brief Add an index to this iterator */ diff --git a/test/benchmarks/transform/CMakeLists.txt b/test/benchmarks/transform/CMakeLists.txt index 5693714..47ee676 100644 --- a/test/benchmarks/transform/CMakeLists.txt +++ b/test/benchmarks/transform/CMakeLists.txt @@ -44,3 +44,21 @@ if(VIKUNJA_ENABLE_CUDA_THRUST_BENCHMARKS) add_test(NAME ${_TARGET_NAME_THRUST_TRANSFORM} COMMAND ${_TARGET_NAME_THRUST_TRANSFORM} ${_VIKUNJA_TEST_OPTIONS}) set_tests_properties(${_TARGET_NAME_THRUST_TRANSFORM} PROPERTIES RUN_SERIAL TRUE) endif() + +set(_TARGET_NAME_VIKUNJA_ZIPITER_TRANSFORM "bench_vikunja_transform_zipiter") + +alpaka_add_executable( + ${_TARGET_NAME_VIKUNJA_ZIPITER_TRANSFORM} + bench_vikunja_transform_zipiter.cpp +) + +target_link_libraries(${_TARGET_NAME_VIKUNJA_ZIPITER_TRANSFORM} + PRIVATE + vikunja::testSetup + vikunja::benchSetup + vikunja::internalvikunja +) + +add_test(NAME ${_TARGET_NAME_VIKUNJA_ZIPITER_TRANSFORM} COMMAND ${_TARGET_NAME_VIKUNJA_ZIPITER_TRANSFORM} ${_VIKUNJA_TEST_OPTIONS}) +# avoid running the benchmarks in parallel +set_tests_properties(${_TARGET_NAME_VIKUNJA_ZIPITER_TRANSFORM} PROPERTIES RUN_SERIAL TRUE) diff --git a/test/benchmarks/transform/bench_vikunja_transform_zipiter.cpp b/test/benchmarks/transform/bench_vikunja_transform_zipiter.cpp new file mode 100644 index 0000000..aabbf7d --- /dev/null +++ b/test/benchmarks/transform/bench_vikunja_transform_zipiter.cpp @@ -0,0 +1,108 @@ +/* Copyright 2022 Anton Reinhard + * + * This file is part of vikunja. + * + * This Source Code Form is subject to the terms of the Mozilla Public + * License, v. 2.0. If a copy of the MPL was not distributed with this + * file, You can obtain one at http://mozilla.org/MPL/2.0/. + */ + +#include +#include +#include +#include +#include + +#include +#include + +#include + +#include + +template +inline void transform_benchmark(TIdx size) +{ + using Setup = vikunja::test::TestAlpakaSetup< + alpaka::DimInt<1u>, // dim + TIdx, // Idx + alpaka::AccCpuSerial, // host type + alpaka::ExampleDefaultAcc, // device type + alpaka::Blocking // queue type + >; + using Vec = alpaka::Vec; + + using TIteratorTuplePtr = std::tuple; + using TIteratorTupleVal = std::tuple; + using ZipIterator = vikunja::mem::iterator::ZipIterator; + + INFO((vikunja::test::print_acc_info(size))); + + Setup setup; + Vec extent = Vec::all(static_cast(size)); + + auto devMemInputInt = vikunja::bench::allocate_mem_iota( + setup, + extent, + static_cast(1), // first value + static_cast(1) // increment + ); + auto devMemInputDouble = vikunja::bench::allocate_mem_iota( + setup, + extent, + static_cast(1), // first value + static_cast(1) // increment + ); + ZipIterator devMemInputPtrBegin(std::make_tuple(alpaka::getPtrNative(devMemInputInt), alpaka::getPtrNative(devMemInputDouble))); + ZipIterator devMemInputPtrEnd = devMemInputPtrBegin + size; + + auto devMemOutput = alpaka::allocBuf(setup.devAcc, extent); + TIteratorTupleVal* devMemOutputPtrBegin = alpaka::getPtrNative(devMemOutput); + + auto hostMemOutput = alpaka::allocBuf(setup.devHost, extent); + TIteratorTupleVal* hostMemOutputPtrBegin = alpaka::getPtrNative(hostMemOutput); + + auto functor = [] ALPAKA_FN_HOST_ACC(TIteratorTupleVal const t) -> auto { return std::make_tuple(2 * std::get<0>(t), 2 * std::get<1>(t)); }; + + vikunja::transform::deviceTransform( + setup.devAcc, + setup.queueAcc, + devMemInputPtrBegin, + devMemInputPtrEnd, + devMemOutputPtrBegin, + functor); + + alpaka::memcpy(setup.queueAcc, hostMemOutput, devMemOutput, extent); + + for(auto i = static_cast(0); i < size; ++i) + { + REQUIRE(static_cast(2) * static_cast(i + 1) == Approx(std::get<0>(hostMemOutputPtrBegin[i]))); + REQUIRE(static_cast(2) * static_cast(i + 1) == Approx(std::get<1>(hostMemOutputPtrBegin[i]))); + } + + // honeypot to check that the function call in the benchmark block has not been removed by the optimizer + hostMemOutputPtrBegin[0] = std::make_tuple(static_cast(42), static_cast(42)); + + BENCHMARK("transform zipiter vikunja") + { + return vikunja::transform::deviceTransform( + setup.devAcc, + setup.queueAcc, + devMemInputPtrBegin, + devMemInputPtrEnd, + devMemOutputPtrBegin, + functor); + }; + + alpaka::memcpy(setup.queueAcc, hostMemOutput, devMemOutput, extent); + + REQUIRE(static_cast(2) == Approx(std::get<0>(hostMemOutputPtrBegin[0]))); + REQUIRE(static_cast(2) == Approx(std::get<1>(hostMemOutputPtrBegin[0]))); +} + +TEMPLATE_TEST_CASE("bechmark transform zipiter", "[benchmark][transform][vikunja][iterator]", int) +{ + using Idx = std::uint64_t; + + transform_benchmark(GENERATE(100, 100'000, 1'270'000, 2'000'000)); +}