From 90a43d2ca78ea08c0037d7e0c768e7de418d1339 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Tom=C3=A1=C5=A1=20Oberhuber?= <oberhuber.tomas@gmail.com> Date: Wed, 31 Mar 2021 21:29:42 +0200 Subject: [PATCH] Removing Args... from segments reductions. --- src/TNL/Algorithms/Segments/BiEllpack.h | 206 +++++++++--------- src/TNL/Algorithms/Segments/BiEllpack.hpp | 12 +- src/TNL/Algorithms/Segments/BiEllpackView.h | 8 +- src/TNL/Algorithms/Segments/BiEllpackView.hpp | 30 ++- src/TNL/Algorithms/Segments/CSR.h | 8 +- src/TNL/Algorithms/Segments/CSR.hpp | 12 +- src/TNL/Algorithms/Segments/CSRView.h | 8 +- src/TNL/Algorithms/Segments/CSRView.hpp | 14 +- src/TNL/Algorithms/Segments/ChunkedEllpack.h | 8 +- .../Algorithms/Segments/ChunkedEllpack.hpp | 12 +- .../Algorithms/Segments/ChunkedEllpackView.h | 8 +- .../Segments/ChunkedEllpackView.hpp | 32 ++- src/TNL/Algorithms/Segments/Ellpack.h | 8 +- src/TNL/Algorithms/Segments/Ellpack.hpp | 12 +- src/TNL/Algorithms/Segments/EllpackView.h | 8 +- src/TNL/Algorithms/Segments/EllpackView.hpp | 22 +- src/TNL/Algorithms/Segments/SlicedEllpack.h | 8 +- src/TNL/Algorithms/Segments/SlicedEllpack.hpp | 12 +- .../Algorithms/Segments/SlicedEllpackView.h | 8 +- .../Algorithms/Segments/SlicedEllpackView.hpp | 22 +- 20 files changed, 225 insertions(+), 233 deletions(-) diff --git a/src/TNL/Algorithms/Segments/BiEllpack.h b/src/TNL/Algorithms/Segments/BiEllpack.h index 347bfe6286..9f0aee6132 100644 --- a/src/TNL/Algorithms/Segments/BiEllpack.h +++ b/src/TNL/Algorithms/Segments/BiEllpack.h @@ -22,163 +22,163 @@ namespace TNL namespace Segments { - template <typename Device, - typename Index, - typename IndexAllocator = typename Allocators::Default<Device>::template Allocator<Index>, - ElementsOrganization Organization = Algorithms::Segments::DefaultElementsOrganization<Device>::getOrganization(), - int WarpSize = 32> - class BiEllpack - { - public: - using DeviceType = Device; - using IndexType = std::remove_const_t<Index>; - using OffsetsHolder = Containers::Vector< IndexType, DeviceType, IndexType, IndexAllocator>; - static constexpr ElementsOrganization getOrganization() { return Organization; } - using ViewType = BiEllpackView< Device, Index, Organization, WarpSize >; - template <typename Device_, typename Index_> - using ViewTemplate = BiEllpackView<Device_, Index_, Organization, WarpSize >; - using ConstViewType = typename ViewType::ConstViewType; - using SegmentViewType = typename ViewType::SegmentViewType; +template <typename Device, + typename Index, + typename IndexAllocator = typename Allocators::Default<Device>::template Allocator<Index>, + ElementsOrganization Organization = Algorithms::Segments::DefaultElementsOrganization<Device>::getOrganization(), + int WarpSize = 32> +class BiEllpack +{ + public: + using DeviceType = Device; + using IndexType = std::remove_const_t<Index>; + using OffsetsHolder = Containers::Vector< IndexType, DeviceType, IndexType, IndexAllocator>; + static constexpr ElementsOrganization getOrganization() { return Organization; } + using ViewType = BiEllpackView< Device, Index, Organization, WarpSize >; + template <typename Device_, typename Index_> + using ViewTemplate = BiEllpackView<Device_, Index_, Organization, WarpSize >; + using ConstViewType = typename ViewType::ConstViewType; + using SegmentViewType = typename ViewType::SegmentViewType; - static constexpr bool havePadding() { return true; }; + static constexpr bool havePadding() { return true; }; - BiEllpack() = default; + BiEllpack() = default; - BiEllpack(const Containers::Vector<IndexType, DeviceType, IndexType> &sizes); + BiEllpack(const Containers::Vector<IndexType, DeviceType, IndexType> &sizes); - BiEllpack(const BiEllpack &segments); + BiEllpack(const BiEllpack &segments); - BiEllpack(const BiEllpack &&segments); + BiEllpack(const BiEllpack &&segments); - static String getSerializationType(); + static String getSerializationType(); - static String getSegmentsType(); + static String getSegmentsType(); - ViewType getView(); + ViewType getView(); - const ConstViewType getConstView() const; + const ConstViewType getConstView() const; - /** + /** * \brief Number of segments. */ - __cuda_callable__ - IndexType - getSegmentsCount() const; + __cuda_callable__ + IndexType + getSegmentsCount() const; - /** + /** * \brief Set sizes of particular segments. */ - template <typename SizesHolder = OffsetsHolder> - void setSegmentsSizes(const SizesHolder &sizes); + template <typename SizesHolder = OffsetsHolder> + void setSegmentsSizes(const SizesHolder &sizes); - void reset(); + void reset(); - IndexType getSegmentSize(const IndexType segmentIdx) const; + IndexType getSegmentSize(const IndexType segmentIdx) const; - /** + /** * \brief Number segments. */ - __cuda_callable__ - IndexType - getSize() const; + __cuda_callable__ + IndexType + getSize() const; - __cuda_callable__ - IndexType - getStorageSize() const; + __cuda_callable__ + IndexType + getStorageSize() const; - __cuda_callable__ - IndexType - getGlobalIndex(const IndexType segmentIdx, const IndexType localIdx) const; + __cuda_callable__ + IndexType + getGlobalIndex(const IndexType segmentIdx, const IndexType localIdx) const; - __cuda_callable__ - SegmentViewType - getSegmentView(const IndexType segmentIdx) const; + __cuda_callable__ + SegmentViewType + getSegmentView(const IndexType segmentIdx) const; - /*** + /*** * \brief Go over all segments and for each segment element call * function 'f' with arguments 'args'. The return type of 'f' is bool. * When its true, the for-loop continues. Once 'f' returns false, the for-loop * is terminated. */ - template <typename Function> - void forElements(IndexType first, IndexType last, Function &&f) const; + template< typename Function > + void forElements( IndexType first, IndexType last, Function&& f ) const; - template <typename Function> - void forAllElements(Function &&f) const; + template <typename Function> + void forAllElements(Function&& f ) const; - template <typename Function> - void forSegments(IndexType begin, IndexType end, Function &&f) const; + template <typename Function> + void forSegments(IndexType begin, IndexType end, Function&& f ) const; - template <typename Function> - void forAllSegments(Function &&f) const; + template <typename Function> + void forAllSegments( Function&& f ) const; - /*** + /*** * \brief Go over all segments and perform a reduction in each of them. */ - template <typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args> - void reduceSegments(IndexType first, IndexType last, Fetch &fetch, const Reduction &reduction, ResultKeeper &keeper, const Real &zero, Args... args) const; - - template <typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args> - void reduceAllSegments(Fetch &fetch, const Reduction &reduction, ResultKeeper &keeper, const Real &zero, Args... args) const; + template <typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceSegments(IndexType first, IndexType last, Fetch &fetch, const Reduction &reduction, ResultKeeper &keeper, const Real &zero ) const; - BiEllpack &operator=(const BiEllpack &source) = default; + template <typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceAllSegments(Fetch &fetch, const Reduction &reduction, ResultKeeper &keeper, const Real &zero ) const; - template <typename Device_, typename Index_, typename IndexAllocator_, ElementsOrganization Organization_> - BiEllpack &operator=(const BiEllpack<Device_, Index_, IndexAllocator_, Organization_, WarpSize> &source); + BiEllpack &operator=(const BiEllpack &source) = default; - void save(File &file) const; + template <typename Device_, typename Index_, typename IndexAllocator_, ElementsOrganization Organization_> + BiEllpack &operator=(const BiEllpack<Device_, Index_, IndexAllocator_, Organization_, WarpSize> &source); - void load(File &file); + void save(File &file) const; - void printStructure(std::ostream &str) const; + void load(File &file); - // TODO: nvcc needs this public because of lambda function used inside - template <typename SizesHolder = OffsetsHolder> - void performRowBubbleSort(const SizesHolder &segmentsSize); + void printStructure(std::ostream &str) const; - // TODO: the same as above - template <typename SizesHolder = OffsetsHolder> - void computeColumnSizes(const SizesHolder &segmentsSizes); + // TODO: nvcc needs this public because of lambda function used inside + template <typename SizesHolder = OffsetsHolder> + void performRowBubbleSort(const SizesHolder &segmentsSize); - protected: - static constexpr int getWarpSize() { return WarpSize; }; + // TODO: the same as above + template <typename SizesHolder = OffsetsHolder> + void computeColumnSizes(const SizesHolder &segmentsSizes); - static constexpr int getLogWarpSize() { return std::log2(WarpSize); }; + protected: + static constexpr int getWarpSize() { return WarpSize; }; - template <typename SizesHolder = OffsetsHolder> - void verifyRowPerm(const SizesHolder &segmentsSizes); + static constexpr int getLogWarpSize() { return std::log2(WarpSize); }; - template <typename SizesHolder = OffsetsHolder> - void verifyRowLengths(const SizesHolder &segmentsSizes); + template <typename SizesHolder = OffsetsHolder> + void verifyRowPerm(const SizesHolder &segmentsSizes); - IndexType getStripLength(const IndexType stripIdx) const; + template <typename SizesHolder = OffsetsHolder> + void verifyRowLengths(const SizesHolder &segmentsSizes); - IndexType getGroupLength(const IndexType strip, const IndexType group) const; + IndexType getStripLength(const IndexType stripIdx) const; - IndexType size = 0, storageSize = 0; + IndexType getGroupLength(const IndexType strip, const IndexType group) const; - IndexType virtualRows = 0; + IndexType size = 0, storageSize = 0; - OffsetsHolder rowPermArray; + IndexType virtualRows = 0; - OffsetsHolder groupPointers; + OffsetsHolder rowPermArray; - // TODO: Replace later - __cuda_callable__ Index power(const IndexType number, const IndexType exponent) const - { - if (exponent >= 0) - { - IndexType result = 1; - for (IndexType i = 0; i < exponent; i++) - result *= number; - return result; - } - return 0; - }; + OffsetsHolder groupPointers; - template <typename Device_, typename Index_, typename IndexAllocator_, ElementsOrganization Organization_, int WarpSize_> - friend class BiEllpack; - }; + // TODO: Replace later + __cuda_callable__ Index power(const IndexType number, const IndexType exponent) const + { + if (exponent >= 0) + { + IndexType result = 1; + for (IndexType i = 0; i < exponent; i++) + result *= number; + return result; + } + return 0; + }; + + template <typename Device_, typename Index_, typename IndexAllocator_, ElementsOrganization Organization_, int WarpSize_> + friend class BiEllpack; +}; } // namespace Segments } // namespace Algorithms diff --git a/src/TNL/Algorithms/Segments/BiEllpack.hpp b/src/TNL/Algorithms/Segments/BiEllpack.hpp index 6dfaa0ab13..16de5a45c7 100644 --- a/src/TNL/Algorithms/Segments/BiEllpack.hpp +++ b/src/TNL/Algorithms/Segments/BiEllpack.hpp @@ -497,12 +497,12 @@ template< typename Device, typename IndexAllocator, ElementsOrganization Organization, int WarpSize > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void BiEllpack< Device, Index, IndexAllocator, Organization, WarpSize >:: -reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->getConstView().reduceSegments( first, last, fetch, reduction, keeper, zero, args... ); + this->getConstView().reduceSegments( first, last, fetch, reduction, keeper, zero ); } template< typename Device, @@ -510,12 +510,12 @@ template< typename Device, typename IndexAllocator, ElementsOrganization Organization, int WarpSize > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void BiEllpack< Device, Index, IndexAllocator, Organization, WarpSize >:: -reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero, args... ); + this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero ); } template< typename Device, diff --git a/src/TNL/Algorithms/Segments/BiEllpackView.h b/src/TNL/Algorithms/Segments/BiEllpackView.h index 344f96c00b..f728eb1acb 100644 --- a/src/TNL/Algorithms/Segments/BiEllpackView.h +++ b/src/TNL/Algorithms/Segments/BiEllpackView.h @@ -126,11 +126,11 @@ class BiEllpackView /*** * \brief Go over all segments and perform a reduction in each of them. */ - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; BiEllpackView& operator=( const BiEllpackView& view ); diff --git a/src/TNL/Algorithms/Segments/BiEllpackView.hpp b/src/TNL/Algorithms/Segments/BiEllpackView.hpp index 29d554e0ae..98c5c05c83 100644 --- a/src/TNL/Algorithms/Segments/BiEllpackView.hpp +++ b/src/TNL/Algorithms/Segments/BiEllpackView.hpp @@ -352,10 +352,10 @@ template< typename Device, typename Index, ElementsOrganization Organization, int WarpSize > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void BiEllpackView< Device, Index, Organization, WarpSize >:: -reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { using RealType = typename detail::FetchLambdaAdapter< Index, Fetch >::ReturnType; if( this->getStorageSize() == 0 ) @@ -425,9 +425,9 @@ reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& dim3 cudaGridSize = Cuda::getMaxGridSize(); if( gridIdx == cudaGrids - 1 ) cudaGridSize.x = cudaBlocks % Cuda::getMaxGridSize(); - details::BiEllpackreduceSegmentsKernel< ViewType, IndexType, Fetch, Reduction, ResultKeeper, Real, BlockDim, Args... > + details::BiEllpackreduceSegmentsKernel< ViewType, IndexType, Fetch, Reduction, ResultKeeper, Real, BlockDim > <<< cudaGridSize, cudaBlockSize, sharedMemory >>> - ( *this, gridIdx, first, last, fetch, reduction, keeper, zero, args... ); + ( *this, gridIdx, first, last, fetch, reduction, keeper, zero ); cudaThreadSynchronize(); TNL_CHECK_CUDA_DEVICE; } @@ -439,12 +439,12 @@ template< typename Device, typename Index, ElementsOrganization Organization, int WarpSize > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void BiEllpackView< Device, Index, Organization, WarpSize >:: -reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero, args... ); + this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero ); } template< typename Device, @@ -513,8 +513,7 @@ template< typename Device, typename Reduction, typename ResultKeeper, typename Real, - int BlockDim, - typename... Args > + int BlockDim > __device__ void BiEllpackView< Device, Index, Organization, WarpSize >:: @@ -524,10 +523,9 @@ reduceSegmentsKernelWithAllParameters( IndexType gridIdx, Fetch fetch, Reduction reduction, ResultKeeper keeper, - Real zero, - Args... args ) const + Real zero ) const { - using RealType = decltype( fetch( IndexType(), IndexType(), IndexType(), std::declval< bool& >(), args... ) ); + using RealType = decltype( fetch( IndexType(), IndexType(), IndexType(), std::declval< bool& >() ) ); const IndexType segmentIdx = ( gridIdx * Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x + first; if( segmentIdx >= last ) return; @@ -569,8 +567,7 @@ template< typename Device, typename Reduction, typename ResultKeeper, typename Real, - int BlockDim, - typename... Args > + int BlockDim > __device__ void BiEllpackView< Device, Index, Organization, WarpSize >:: @@ -580,10 +577,9 @@ reduceSegmentsKernel( IndexType gridIdx, Fetch fetch, Reduction reduction, ResultKeeper keeper, - Real zero, - Args... args ) const + Real zero ) const { - using RealType = decltype( fetch( IndexType(), std::declval< bool& >(), args... ) ); + using RealType = decltype( fetch( IndexType(), std::declval< bool& >() ) ); Index segmentIdx = ( gridIdx * Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x + first; const IndexType strip = segmentIdx >> getLogWarpSize(); diff --git a/src/TNL/Algorithms/Segments/CSR.h b/src/TNL/Algorithms/Segments/CSR.h index aee35bbded..f3761ace55 100644 --- a/src/TNL/Algorithms/Segments/CSR.h +++ b/src/TNL/Algorithms/Segments/CSR.h @@ -123,11 +123,11 @@ class CSR /*** * \brief Go over all segments and perform a reduction in each of them. */ - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; CSR& operator=( const CSR& rhsSegments ) = default; diff --git a/src/TNL/Algorithms/Segments/CSR.hpp b/src/TNL/Algorithms/Segments/CSR.hpp index f0e1d78811..de3b776933 100644 --- a/src/TNL/Algorithms/Segments/CSR.hpp +++ b/src/TNL/Algorithms/Segments/CSR.hpp @@ -275,24 +275,24 @@ template< typename Device, typename Index, typename Kernel, typename IndexAllocator > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void CSR< Device, Index, Kernel, IndexAllocator >:: -reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->getConstView().reduceSegments( first, last, fetch, reduction, keeper, zero, args... ); + this->getConstView().reduceSegments( first, last, fetch, reduction, keeper, zero ); } template< typename Device, typename Index, typename Kernel, typename IndexAllocator > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void CSR< Device, Index, Kernel, IndexAllocator >:: -reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero, args... ); + this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero ); } template< typename Device, diff --git a/src/TNL/Algorithms/Segments/CSRView.h b/src/TNL/Algorithms/Segments/CSRView.h index 6b47a4a3a5..39a97b8c55 100644 --- a/src/TNL/Algorithms/Segments/CSRView.h +++ b/src/TNL/Algorithms/Segments/CSRView.h @@ -120,11 +120,11 @@ class CSRView /*** * \brief Go over all segments and perform a reduction in each of them. */ - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; CSRView& operator=( const CSRView& view ); diff --git a/src/TNL/Algorithms/Segments/CSRView.hpp b/src/TNL/Algorithms/Segments/CSRView.hpp index e6cbcfff44..f4cfc2c786 100644 --- a/src/TNL/Algorithms/Segments/CSRView.hpp +++ b/src/TNL/Algorithms/Segments/CSRView.hpp @@ -241,26 +241,26 @@ forAllSegments( Function&& f ) const template< typename Device, typename Index, typename Kernel > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void CSRView< Device, Index, Kernel >:: -reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { if( std::is_same< DeviceType, TNL::Devices::Host >::value ) - TNL::Algorithms::Segments::CSRScalarKernel< IndexType, DeviceType >::reduceSegments( offsets, first, last, fetch, reduction, keeper, zero, args... ); + TNL::Algorithms::Segments::CSRScalarKernel< IndexType, DeviceType >::reduceSegments( offsets, first, last, fetch, reduction, keeper, zero ); else - kernel.reduceSegments( offsets, first, last, fetch, reduction, keeper, zero, args... ); + kernel.reduceSegments( offsets, first, last, fetch, reduction, keeper, zero ); } template< typename Device, typename Index, typename Kernel > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void CSRView< Device, Index, Kernel >:: -reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero, args... ); + this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero ); } template< typename Device, diff --git a/src/TNL/Algorithms/Segments/ChunkedEllpack.h b/src/TNL/Algorithms/Segments/ChunkedEllpack.h index 52ab9cd212..5ed48a7f76 100644 --- a/src/TNL/Algorithms/Segments/ChunkedEllpack.h +++ b/src/TNL/Algorithms/Segments/ChunkedEllpack.h @@ -111,11 +111,11 @@ class ChunkedEllpack /*** * \brief Go over all segments and perform a reduction in each of them. */ - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; ChunkedEllpack& operator=( const ChunkedEllpack& source ) = default; diff --git a/src/TNL/Algorithms/Segments/ChunkedEllpack.hpp b/src/TNL/Algorithms/Segments/ChunkedEllpack.hpp index 3031c7baab..2228b7e620 100644 --- a/src/TNL/Algorithms/Segments/ChunkedEllpack.hpp +++ b/src/TNL/Algorithms/Segments/ChunkedEllpack.hpp @@ -440,24 +440,24 @@ template< typename Device, typename Index, typename IndexAllocator, ElementsOrganization Organization > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void ChunkedEllpack< Device, Index, IndexAllocator, Organization >:: -reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->getConstView().reduceSegments( first, last, fetch, reduction, keeper, zero, args... ); + this->getConstView().reduceSegments( first, last, fetch, reduction, keeper, zero ); } template< typename Device, typename Index, typename IndexAllocator, ElementsOrganization Organization > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void ChunkedEllpack< Device, Index, IndexAllocator, Organization >:: -reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero, args... ); + this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero ); } template< typename Device, diff --git a/src/TNL/Algorithms/Segments/ChunkedEllpackView.h b/src/TNL/Algorithms/Segments/ChunkedEllpackView.h index 77352f0815..8f8177dbb4 100644 --- a/src/TNL/Algorithms/Segments/ChunkedEllpackView.h +++ b/src/TNL/Algorithms/Segments/ChunkedEllpackView.h @@ -140,11 +140,11 @@ class ChunkedEllpackView /*** * \brief Go over all segments and perform a reduction in each of them. */ - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; ChunkedEllpackView& operator=( const ChunkedEllpackView& view ); diff --git a/src/TNL/Algorithms/Segments/ChunkedEllpackView.hpp b/src/TNL/Algorithms/Segments/ChunkedEllpackView.hpp index 92bc43b775..c43b57321d 100644 --- a/src/TNL/Algorithms/Segments/ChunkedEllpackView.hpp +++ b/src/TNL/Algorithms/Segments/ChunkedEllpackView.hpp @@ -393,15 +393,15 @@ forAllSegments( Function&& f ) const template< typename Device, typename Index, ElementsOrganization Organization > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void ChunkedEllpackView< Device, Index, Organization >:: -reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { using RealType = typename detail::FetchLambdaAdapter< Index, Fetch >::ReturnType; if( std::is_same< DeviceType, Devices::Host >::value ) { - //reduceSegmentsKernel( 0, first, last, fetch, reduction, keeper, zero, args... ); + //reduceSegmentsKernel( 0, first, last, fetch, reduction, keeper, zero ); //return; for( IndexType segmentIdx = first; segmentIdx < last; segmentIdx++ ) @@ -456,9 +456,9 @@ reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& { if( gridIdx == cudaGrids - 1 ) cudaGridSize.x = cudaBlocks % Cuda::getMaxGridSize(); - details::ChunkedEllpackreduceSegmentsKernel< ViewType, IndexType, Fetch, Reduction, ResultKeeper, Real, Args... > + details::ChunkedEllpackreduceSegmentsKernel< ViewType, IndexType, Fetch, Reduction, ResultKeeper, Real > <<< cudaGridSize, cudaBlockSize, sharedMemory >>> - ( *this, gridIdx, first, last, fetch, reduction, keeper, zero, args... ); + ( *this, gridIdx, first, last, fetch, reduction, keeper, zero ); } #endif } @@ -467,12 +467,12 @@ reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& template< typename Device, typename Index, ElementsOrganization Organization > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void ChunkedEllpackView< Device, Index, Organization >:: -reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero, args... ); + this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero ); } template< typename Device, @@ -543,8 +543,7 @@ template< typename Device, template< typename Fetch, typename Reduction, typename ResultKeeper, - typename Real, - typename... Args > + typename Real > __device__ void ChunkedEllpackView< Device, Index, Organization >:: @@ -554,10 +553,9 @@ reduceSegmentsKernelWithAllParameters( IndexType gridIdx, Fetch fetch, Reduction reduction, ResultKeeper keeper, - Real zero, - Args... args ) const + Real zero ) const { - using RealType = decltype( fetch( IndexType(), IndexType(), IndexType(), std::declval< bool& >(), args... ) ); + using RealType = decltype( fetch( IndexType(), IndexType(), IndexType(), std::declval< bool& >() ) ); const IndexType firstSlice = rowToSliceMapping[ first ]; const IndexType lastSlice = rowToSliceMapping[ last - 1 ]; @@ -621,8 +619,7 @@ template< typename Device, template< typename Fetch, typename Reduction, typename ResultKeeper, - typename Real, - typename... Args > + typename Real > __device__ void ChunkedEllpackView< Device, Index, Organization >:: @@ -632,10 +629,9 @@ reduceSegmentsKernel( IndexType gridIdx, Fetch fetch, Reduction reduction, ResultKeeper keeper, - Real zero, - Args... args ) const + Real zero ) const { - using RealType = decltype( fetch( IndexType(), std::declval< bool& >(), args... ) ); + using RealType = decltype( fetch( IndexType(), std::declval< bool& >() ) ); const IndexType firstSlice = rowToSliceMapping[ first ]; const IndexType lastSlice = rowToSliceMapping[ last - 1 ]; diff --git a/src/TNL/Algorithms/Segments/Ellpack.h b/src/TNL/Algorithms/Segments/Ellpack.h index 368d435eb0..52372e64b9 100644 --- a/src/TNL/Algorithms/Segments/Ellpack.h +++ b/src/TNL/Algorithms/Segments/Ellpack.h @@ -111,11 +111,11 @@ class Ellpack /*** * \brief Go over all segments and perform a reduction in each of them. */ - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; Ellpack& operator=( const Ellpack& source ) = default; diff --git a/src/TNL/Algorithms/Segments/Ellpack.hpp b/src/TNL/Algorithms/Segments/Ellpack.hpp index 0355ee62b0..9de1f5b970 100644 --- a/src/TNL/Algorithms/Segments/Ellpack.hpp +++ b/src/TNL/Algorithms/Segments/Ellpack.hpp @@ -304,12 +304,12 @@ template< typename Device, typename IndexAllocator, ElementsOrganization Organization, int Alignment > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void Ellpack< Device, Index, IndexAllocator, Organization, Alignment >:: -reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->getConstView().reduceSegments( first, last, fetch, reduction, keeper, zero, args... ); + this->getConstView().reduceSegments( first, last, fetch, reduction, keeper, zero ); } template< typename Device, @@ -317,12 +317,12 @@ template< typename Device, typename IndexAllocator, ElementsOrganization Organization, int Alignment > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void Ellpack< Device, Index, IndexAllocator, Organization, Alignment >:: -reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero, args... ); + this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero ); } template< typename Device, diff --git a/src/TNL/Algorithms/Segments/EllpackView.h b/src/TNL/Algorithms/Segments/EllpackView.h index 37e6290dbc..23df8813b2 100644 --- a/src/TNL/Algorithms/Segments/EllpackView.h +++ b/src/TNL/Algorithms/Segments/EllpackView.h @@ -110,11 +110,11 @@ class EllpackView /*** * \brief Go over all segments and perform a reduction in each of them. */ - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; EllpackView& operator=( const EllpackView& view ); diff --git a/src/TNL/Algorithms/Segments/EllpackView.hpp b/src/TNL/Algorithms/Segments/EllpackView.hpp index b399151855..120f8f4369 100644 --- a/src/TNL/Algorithms/Segments/EllpackView.hpp +++ b/src/TNL/Algorithms/Segments/EllpackView.hpp @@ -271,16 +271,16 @@ template< typename Device, typename Index, ElementsOrganization Organization, int Alignment > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void EllpackView< Device, Index, Organization, Alignment >:: -reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - //using RealType = decltype( fetch( IndexType(), IndexType(), IndexType(), std::declval< bool& >(), args... ) ); - using RealType = typename detail::FetchLambdaAdapter< Index, Fetch >::ReturnType; + //using RealType = decltype( fetch( IndexType(), IndexType(), IndexType(), std::declval< bool& >() ) ); + using RealType = typename details::FetchLambdaAdapter< Index, Fetch >::ReturnType; if( Organization == RowMajorOrder ) { const IndexType segmentSize = this->segmentSize; - auto l = [=] __cuda_callable__ ( const IndexType segmentIdx, Args... args ) mutable { + auto l = [=] __cuda_callable__ ( const IndexType segmentIdx ) mutable { const IndexType begin = segmentIdx * segmentSize; const IndexType end = begin + segmentSize; RealType aux( zero ); @@ -290,13 +290,13 @@ reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& aux = reduction( aux, detail::FetchLambdaAdapter< IndexType, Fetch >::call( fetch, segmentIdx, localIdx++, j, compute ) ); keeper( segmentIdx, aux ); }; - Algorithms::ParallelFor< Device >::exec( first, last, l, args... ); + Algorithms::ParallelFor< Device >::exec( first, last, l ); } else { const IndexType storageSize = this->getStorageSize(); const IndexType alignedSize = this->alignedSize; - auto l = [=] __cuda_callable__ ( const IndexType segmentIdx, Args... args ) mutable { + auto l = [=] __cuda_callable__ ( const IndexType segmentIdx ) mutable { const IndexType begin = segmentIdx; const IndexType end = storageSize; RealType aux( zero ); @@ -306,7 +306,7 @@ reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& aux = reduction( aux, detail::FetchLambdaAdapter< IndexType, Fetch >::call( fetch, segmentIdx, localIdx++, j, compute ) ); keeper( segmentIdx, aux ); }; - Algorithms::ParallelFor< Device >::exec( first, last, l, args... ); + Algorithms::ParallelFor< Device >::exec( first, last, l ); } } @@ -314,11 +314,11 @@ template< typename Device, typename Index, ElementsOrganization Organization, int Alignment > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void EllpackView< Device, Index, Organization, Alignment >:: -reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero, args... ); + this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero ); } template< typename Device, diff --git a/src/TNL/Algorithms/Segments/SlicedEllpack.h b/src/TNL/Algorithms/Segments/SlicedEllpack.h index edb35fac3b..c69a271420 100644 --- a/src/TNL/Algorithms/Segments/SlicedEllpack.h +++ b/src/TNL/Algorithms/Segments/SlicedEllpack.h @@ -108,11 +108,11 @@ class SlicedEllpack /*** * \brief Go over all segments and perform a reduction in each of them. */ - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; SlicedEllpack& operator=( const SlicedEllpack& source ) = default; diff --git a/src/TNL/Algorithms/Segments/SlicedEllpack.hpp b/src/TNL/Algorithms/Segments/SlicedEllpack.hpp index 8f90b3aa5a..af41847478 100644 --- a/src/TNL/Algorithms/Segments/SlicedEllpack.hpp +++ b/src/TNL/Algorithms/Segments/SlicedEllpack.hpp @@ -338,12 +338,12 @@ template< typename Device, typename IndexAllocator, ElementsOrganization Organization, int SliceSize > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void SlicedEllpack< Device, Index, IndexAllocator, Organization, SliceSize >:: -reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->getConstView().reduceSegments( first, last, fetch, reduction, keeper, zero, args... ); + this->getConstView().reduceSegments( first, last, fetch, reduction, keeper, zero ); } template< typename Device, @@ -351,12 +351,12 @@ template< typename Device, typename IndexAllocator, ElementsOrganization Organization, int SliceSize > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void SlicedEllpack< Device, Index, IndexAllocator, Organization, SliceSize >:: -reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero, args... ); + this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero ); } template< typename Device, diff --git a/src/TNL/Algorithms/Segments/SlicedEllpackView.h b/src/TNL/Algorithms/Segments/SlicedEllpackView.h index 5b052848c4..634d193c08 100644 --- a/src/TNL/Algorithms/Segments/SlicedEllpackView.h +++ b/src/TNL/Algorithms/Segments/SlicedEllpackView.h @@ -109,11 +109,11 @@ class SlicedEllpackView /*** * \brief Go over all segments and perform a reduction in each of them. */ - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > - void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const; + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > + void reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const; SlicedEllpackView& operator=( const SlicedEllpackView& view ); diff --git a/src/TNL/Algorithms/Segments/SlicedEllpackView.hpp b/src/TNL/Algorithms/Segments/SlicedEllpackView.hpp index 13ef6b0388..703da8d3c3 100644 --- a/src/TNL/Algorithms/Segments/SlicedEllpackView.hpp +++ b/src/TNL/Algorithms/Segments/SlicedEllpackView.hpp @@ -326,18 +326,18 @@ template< typename Device, typename Index, ElementsOrganization Organization, int SliceSize > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void SlicedEllpackView< Device, Index, Organization, SliceSize >:: -reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - using RealType = typename detail::FetchLambdaAdapter< Index, Fetch >::ReturnType; - //using RealType = decltype( fetch( IndexType(), IndexType(), IndexType(), std::declval< bool& >(), args... ) ); + using RealType = typename details::FetchLambdaAdapter< Index, Fetch >::ReturnType; + //using RealType = decltype( fetch( IndexType(), IndexType(), IndexType(), std::declval< bool& >() ) ); const auto sliceSegmentSizes_view = this->sliceSegmentSizes.getConstView(); const auto sliceOffsets_view = this->sliceOffsets.getConstView(); if( Organization == RowMajorOrder ) { - auto l = [=] __cuda_callable__ ( const IndexType segmentIdx, Args... args ) mutable { + auto l = [=] __cuda_callable__ ( const IndexType segmentIdx ) mutable { const IndexType sliceIdx = segmentIdx / SliceSize; const IndexType segmentInSliceIdx = segmentIdx % SliceSize; const IndexType segmentSize = sliceSegmentSizes_view[ sliceIdx ]; @@ -350,11 +350,11 @@ reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& aux = reduction( aux, detail::FetchLambdaAdapter< IndexType, Fetch >::call( fetch, segmentIdx, localIdx++, globalIdx, compute ) ); keeper( segmentIdx, aux ); }; - Algorithms::ParallelFor< Device >::exec( first, last, l, args... ); + Algorithms::ParallelFor< Device >::exec( first, last, l ); } else { - auto l = [=] __cuda_callable__ ( const IndexType segmentIdx, Args... args ) mutable { + auto l = [=] __cuda_callable__ ( const IndexType segmentIdx ) mutable { const IndexType sliceIdx = segmentIdx / SliceSize; const IndexType segmentInSliceIdx = segmentIdx % SliceSize; //const IndexType segmentSize = sliceSegmentSizes_view[ sliceIdx ]; @@ -367,7 +367,7 @@ reduceSegments( IndexType first, IndexType last, Fetch& fetch, const Reduction& aux = reduction( aux, detail::FetchLambdaAdapter< IndexType, Fetch >::call( fetch, segmentIdx, localIdx++, globalIdx, compute ) ); keeper( segmentIdx, aux ); }; - Algorithms::ParallelFor< Device >::exec( first, last, l, args... ); + Algorithms::ParallelFor< Device >::exec( first, last, l ); } } @@ -375,12 +375,12 @@ template< typename Device, typename Index, ElementsOrganization Organization, int SliceSize > - template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real, typename... Args > + template< typename Fetch, typename Reduction, typename ResultKeeper, typename Real > void SlicedEllpackView< Device, Index, Organization, SliceSize >:: -reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero, Args... args ) const +reduceAllSegments( Fetch& fetch, const Reduction& reduction, ResultKeeper& keeper, const Real& zero ) const { - this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero, args... ); + this->reduceSegments( 0, this->getSegmentsCount(), fetch, reduction, keeper, zero ); } template< typename Device, -- GitLab