diff --git a/src/TNL/Algorithms/Segments/BiEllpack.h b/src/TNL/Algorithms/Segments/BiEllpack.h index 347bfe6286ded5e88a9d1e42443c2ea77a452ec0..9f0aee61329b51468f34faf4e703d33b38ba7ae7 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 6dfaa0ab13ec284a511bd1ceea3bbbe230bd2f96..16de5a45c7814574337f96825562e8b70fc4de50 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 344f96c00bbe12e8326180a8a30b92519f14a76f..f728eb1acb952458f936749f63e31a1b177bd576 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 29d554e0aee65062ae5f52785da04fa245e46536..98c5c05c83cee68f9f62b4f8b79063be6c7e2c6b 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 aee35bbded6f1d966ebbec910401f19b23d4eaa0..f3761ace552afa5bc452a49d60c3b577c4e345ee 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 f0e1d788116c62e694702eaae168d12cff98832d..de3b7769336d88e37fdf44c6d5570775f438face 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 6b47a4a3a5bfcfa6a6083b63eb4d6130e5e37d10..39a97b8c55cfbc85673f9c69cad6d8b2f6caf6ea 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 e6cbcfff44c21a0d264f88a449885e5b45e21498..f4cfc2c786e37c9d73ecd617d6c1d3591f7a5600 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 52ab9cd212fc67abed78df39bce5e60b10601fab..5ed48a7f7673bccd60aeee750694dfea8f2a8408 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 3031c7baab6dabfe7986bc06dc238b6668733fa4..2228b7e62041e647fecc2eb19892c6ffaf53bd64 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 77352f081525f34df33d411ff794031d912bb5e9..8f8177dbb4309e572dac4010285d706bd3c5dd4b 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 92bc43b775e62f29feffc16aa1c9a3d5c66de0ad..c43b57321d853b00315978776cc31bab6e216838 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 368d435eb0b8322312a5e45e241422349bea12c2..52372e64b9f54ca2b95280bb6d51cd97e9cd0a82 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 0355ee62b0639aedd3dd8886df0e8aa24871f497..9de1f5b9705c7ebe024a78756da3a9211865c8ff 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 37e6290dbc344be8fafb6527c5d4e569a1b85ca0..23df8813b22bef08ecaa46bbf26f1b5c303bbda1 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 b399151855735575aef7a19584ea9d0da49ab42e..120f8f43692d150870b799c04ac52647f8e8583c 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 edb35fac3b660f0ffd828f49033b82a5bfa14241..c69a27142091e78b8dbcde7718c313e6b100ef7e 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 8f90b3aa5a2f3d16fd53c6d527548f5ddadf0f31..af418474781c3dbbabb8149d00d133bb5e0dc0cf 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 5b052848c4c4e8a9ba4c2905ed0e3baa5df447c7..634d193c0856d5ca5fce188f00e0ac2217fae449 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 13ef6b038864560ac89a06d8b8779ce110535304..703da8d3c32e20f2cef99f84fbcbee6f56415327 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,