diff --git a/src/TNL/Algorithms/MemoryOperations.h b/src/TNL/Algorithms/MemoryOperations.h index 68a3a88146a5f854fd0ba4c8852fd5c2d9d60761..42c37f062c62e7ae7108fe0646bbde72d4dd9d15 100644 --- a/src/TNL/Algorithms/MemoryOperations.h +++ b/src/TNL/Algorithms/MemoryOperations.h @@ -24,6 +24,25 @@ struct MemoryOperations; template<> struct MemoryOperations< Devices::Sequential > { + template< typename Element, typename Index > + __cuda_callable__ + static void construct( Element* data, + const Index size ); + + // note that args are passed by reference to the constructor, not via + // std::forward since move-semantics does not apply for the construction of + // multiple elements + template< typename Element, typename Index, typename... Args > + __cuda_callable__ + static void construct( Element* data, + const Index size, + const Args&... args ); + + template< typename Element, typename Index > + __cuda_callable__ + static void destruct( Element* data, + const Index size ); + template< typename Element > __cuda_callable__ static void setElement( Element* data, @@ -81,6 +100,22 @@ struct MemoryOperations< Devices::Sequential > template<> struct MemoryOperations< Devices::Host > { + template< typename Element, typename Index > + static void construct( Element* data, + const Index size ); + + // note that args are passed by reference to the constructor, not via + // std::forward since move-semantics does not apply for the construction of + // multiple elements + template< typename Element, typename Index, typename... Args > + static void construct( Element* data, + const Index size, + const Args&... args ); + + template< typename Element, typename Index > + static void destruct( Element* data, + const Index size ); + // this is __cuda_callable__ only to silence nvcc warnings TNL_NVCC_HD_WARNING_DISABLE template< typename Element > @@ -137,6 +172,23 @@ struct MemoryOperations< Devices::Host > template<> struct MemoryOperations< Devices::Cuda > { + template< typename Element, typename Index > + static void construct( Element* data, + const Index size ); + + // note that args are passed by value to the constructor, not via + // std::forward or even by reference, since move-semantics does not apply for + // the construction of multiple elements and pass-by-reference cannot be used + // with CUDA kernels + template< typename Element, typename Index, typename... Args > + static void construct( Element* data, + const Index size, + const Args&... args ); + + template< typename Element, typename Index > + static void destruct( Element* data, + const Index size ); + template< typename Element > __cuda_callable__ static void setElement( Element* data, diff --git a/src/TNL/Algorithms/MemoryOperationsCuda.hpp b/src/TNL/Algorithms/MemoryOperationsCuda.hpp index 5351b69625e22e73b5039b39321a6338c5944f2f..8c104469b5154beeab3f1ec782346c820c78536c 100644 --- a/src/TNL/Algorithms/MemoryOperationsCuda.hpp +++ b/src/TNL/Algorithms/MemoryOperationsCuda.hpp @@ -23,6 +23,56 @@ namespace TNL { namespace Algorithms { +template< typename Element, typename Index > +void +MemoryOperations< Devices::Cuda >:: +construct( Element* data, + const Index size ) +{ + TNL_ASSERT_TRUE( data, "Attempted to create elements through a nullptr." ); + auto kernel = [data] __cuda_callable__ ( Index i ) + { + // placement-new + ::new( (void*) (data + i) ) Element(); + }; + ParallelFor< Devices::Cuda >::exec( (Index) 0, size, kernel ); +} + +template< typename Element, typename Index, typename... Args > +void +MemoryOperations< Devices::Cuda >:: +construct( Element* data, + const Index size, + const Args&... args ) +{ + TNL_ASSERT_TRUE( data, "Attempted to create elements through a nullptr." ); + // NOTE: nvcc does not allow __cuda_callable__ lambdas with a variadic capture + auto kernel = [data] __cuda_callable__ ( Index i, Args... args ) + { + // placement-new + // (note that args are passed by value to the constructor, not via + // std::forward or even by reference, since move-semantics does not apply for + // the construction of multiple elements and pass-by-reference cannot be used + // with CUDA kernels) + ::new( (void*) (data + i) ) Element( args... ); + }; + ParallelFor< Devices::Cuda >::exec( (Index) 0, size, kernel, args... ); +} + +template< typename Element, typename Index > +void +MemoryOperations< Devices::Cuda >:: +destruct( Element* data, + const Index size ) +{ + TNL_ASSERT_TRUE( data, "Attempted to destroy data through a nullptr." ); + auto kernel = [data] __cuda_callable__ ( Index i ) + { + (data + i)->~Element(); + }; + ParallelFor< Devices::Cuda >::exec( (Index) 0, size, kernel ); +} + template< typename Element > __cuda_callable__ void MemoryOperations< Devices::Cuda >:: @@ -57,7 +107,7 @@ getElement( const Element* data ) return *data; #else // TODO: For some reason the following does not work after adding - // #ifdef __CUDA_ARCH__ to Array::getElement and ArrayView::getElement + // #ifdef __CUDA_ARCH__ to Array::getElement and ArrayView::getElement // Probably it might be a problem with lambda function 'kernel' which // nvcc probably does not handle properly. //MultiDeviceMemoryOperations< void, Devices::Cuda >::template copy< Element, Element, int >( &result, data, 1 ); diff --git a/src/TNL/Algorithms/MemoryOperationsHost.hpp b/src/TNL/Algorithms/MemoryOperationsHost.hpp index 92b44f8cf51fe085eda52d41f98524125631b6d4..0034b8302c98cf1b657b76941f6dd05935161dd0 100644 --- a/src/TNL/Algorithms/MemoryOperationsHost.hpp +++ b/src/TNL/Algorithms/MemoryOperationsHost.hpp @@ -21,6 +21,54 @@ namespace TNL { namespace Algorithms { +template< typename Element, typename Index > +void +MemoryOperations< Devices::Host >:: +construct( Element* data, + const Index size ) +{ + TNL_ASSERT_TRUE( data, "Attempted to create elements through a nullptr." ); + auto kernel = [data]( Index i ) + { + // placement-new + ::new( (void*) (data + i) ) Element(); + }; + ParallelFor< Devices::Host >::exec( (Index) 0, size, kernel ); +} + +template< typename Element, typename Index, typename... Args > +void +MemoryOperations< Devices::Host >:: +construct( Element* data, + const Index size, + const Args&... args ) +{ + TNL_ASSERT_TRUE( data, "Attempted to create elements through a nullptr." ); + auto kernel = [data, &args...]( Index i ) + { + // placement-new + // (note that args are passed by reference to the constructor, not via + // std::forward since move-semantics does not apply for the construction + // of multiple elements) + ::new( (void*) (data + i) ) Element( args... ); + }; + ParallelFor< Devices::Host >::exec( (Index) 0, size, kernel ); +} + +template< typename Element, typename Index > +void +MemoryOperations< Devices::Host >:: +destruct( Element* data, + const Index size ) +{ + TNL_ASSERT_TRUE( data, "Attempted to destroy data through a nullptr." ); + auto kernel = [data]( Index i ) + { + (data + i)->~Element(); + }; + ParallelFor< Devices::Host >::exec( (Index) 0, size, kernel ); +} + template< typename Element > __cuda_callable__ // only to avoid nvcc warning void diff --git a/src/TNL/Algorithms/MemoryOperationsSequential.hpp b/src/TNL/Algorithms/MemoryOperationsSequential.hpp index 9e5ad25b1392ccd093952b0dbb25b941370eb833..2ea21d0aca0ead7ab3bcf7d6b95a96599d1f27e3 100644 --- a/src/TNL/Algorithms/MemoryOperationsSequential.hpp +++ b/src/TNL/Algorithms/MemoryOperationsSequential.hpp @@ -15,6 +15,48 @@ namespace TNL { namespace Algorithms { +template< typename Element, typename Index > +__cuda_callable__ +void +MemoryOperations< Devices::Sequential >:: +construct( Element* data, + const Index size ) +{ + TNL_ASSERT_TRUE( data, "Attempted to create elements through a nullptr." ); + for( Index i = 0; i < size; i++ ) + // placement-new + ::new( (void*) (data + i) ) Element(); +} + +template< typename Element, typename Index, typename... Args > +__cuda_callable__ +void +MemoryOperations< Devices::Sequential >:: +construct( Element* data, + const Index size, + const Args&... args ) +{ + TNL_ASSERT_TRUE( data, "Attempted to create elements through a nullptr." ); + for( Index i = 0; i < size; i++ ) + // placement-new + // (note that args are passed by reference to the constructor, not via + // std::forward since move-semantics does not apply for the construction + // of multiple elements) + ::new( (void*) (data + i) ) Element( args... ); +} + +template< typename Element, typename Index > +__cuda_callable__ +void +MemoryOperations< Devices::Sequential >:: +destruct( Element* data, + const Index size ) +{ + TNL_ASSERT_TRUE( data, "Attempted to destroy elements through a nullptr." ); + for( Index i = 0; i < size; i++ ) + (data + i)->~Element(); +} + template< typename Element > __cuda_callable__ void @@ -22,6 +64,7 @@ MemoryOperations< Devices::Sequential >:: setElement( Element* data, const Element& value ) { + TNL_ASSERT_TRUE( data, "Attempted to set data through a nullptr." ); *data = value; } @@ -31,6 +74,7 @@ Element MemoryOperations< Devices::Sequential >:: getElement( const Element* data ) { + TNL_ASSERT_TRUE( data, "Attempted to get data through a nullptr." ); return *data; } @@ -42,7 +86,9 @@ set( Element* data, const Element& value, const Index size ) { - for( Index i = 0; i < size; i ++ ) + if( size == 0 ) return; + TNL_ASSERT_TRUE( data, "Attempted to set data through a nullptr." ); + for( Index i = 0; i < size; i++ ) data[ i ] = value; } @@ -56,7 +102,11 @@ copy( DestinationElement* destination, const SourceElement* source, const Index size ) { - for( Index i = 0; i < size; i ++ ) + if( size == 0 ) return; + TNL_ASSERT_TRUE( destination, "Attempted to copy data to a nullptr." ); + TNL_ASSERT_TRUE( source, "Attempted to copy data from a nullptr." ); + + for( Index i = 0; i < size; i++ ) destination[ i ] = source[ i ]; } @@ -87,6 +137,10 @@ compare( const Element1* destination, const Element2* source, const Index size ) { + if( size == 0 ) return true; + TNL_ASSERT_TRUE( destination, "Attempted to compare data through a nullptr." ); + TNL_ASSERT_TRUE( source, "Attempted to compare data through a nullptr." ); + for( Index i = 0; i < size; i++ ) if( ! ( destination[ i ] == source[ i ] ) ) return false; diff --git a/src/TNL/Containers/Array.h b/src/TNL/Containers/Array.h index 92a976e7262cd8b9a73957f67e8a99a849166786..9c02ac9d46c028d197f935693bfc8502cec8ad9b 100644 --- a/src/TNL/Containers/Array.h +++ b/src/TNL/Containers/Array.h @@ -41,9 +41,9 @@ template< int, typename > class StaticArray; * * Memory management handled by constructors and destructors according to the * [RAII](https://en.wikipedia.org/wiki/RAII) principle and by methods - * \ref setSize, \ref setLike, \ref swap, and \ref reset. You can also use - * methods \ref getSize and \ref empty to check the current array size and - * \ref getData to access the raw pointer. + * \ref resize \ref setSize, \ref setLike, \ref swap, and \ref reset. You can + * also use methods \ref getSize and \ref empty to check the current array size + * and \ref getData to access the raw pointer. * * Methods annotated as \ref \_\_cuda_callable\_\_ can be called either from * host or from kernels executing on a device according to the \e Device @@ -263,16 +263,60 @@ class Array virtual String getSerializationTypeVirtual() const; /** - * \brief Method for setting the array size. + * \brief Method for resizing the array. + * + * The method resizes the array to the given size: + * + * - If the current size is greater than count, the array is reduced to + * its first \e size elements. + * - If the current size is less than \e size, additional elements are + * appended (see the note below on initialization). + * - If the current size is equal to \e size, nothing happens. + * + * If the array size changes, the current data will be deallocated, thus + * all pointers and views to the array alements will become invalid. + * + * Note that this method differs from \ref std::vector::resize with respect + * to the initialization of array elements: + * + * - if \e ValueType is a [fundamental type](https://en.cppreference.com/w/cpp/types/is_fundamental), + * the elements are [default-initialized](https://en.cppreference.com/w/cpp/language/default_initialization) + * (i.e., the elements are initialized to indeterminate values). + * - otherwise, the elements are [value-initialized](https://en.cppreference.com/w/cpp/language/value_initialization) + * (like in \ref std::vector::resize). + * + * \param size The new size of the array. + */ + void resize( Index size ); + + /** + * \brief Method for resizing the array with an initial value. + * + * The method resizes the array to the given size: * - * If the array shares data with other arrays, the data is unbound. If the - * current data is not shared and the current size is the same as the new - * one, nothing happens. + * - If the current size is greater than count, the array is reduced to + * its first \e size elements. + * - If the current size is less than \e size, additional copies of + * \e value are appended. + * - If the current size is equal to \e size, nothing happens. * * If the array size changes, the current data will be deallocated, thus * all pointers and views to the array alements will become invalid. * * \param size The new size of the array. + * \param value The value to initialize new elements with. + */ + void resize( Index size, const ValueType& value ); + + /** + * \brief Method for setting the array size. + * + * This method behaves almost like \ref resize, but when the array size + * is changed, old elements are not copied to the new memory location. + * Hence, this is a shortcut for deallocating the array with \e resize(0) + * followed by setting the new size with \e resize(size) + * + * \param size The new size of the array. */ void setSize( Index size ); @@ -289,6 +333,8 @@ class Array * If the array size changes, the current data will be deallocated, thus * all pointers and views to the array alements will become invalid. * + * Note that this method uses \ref setSize rather than \ref resize. + * * \tparam ArrayT The type of the parameter can be any type which provides * the method \ref getSize() with the same signature as \e Array. * \param array The array whose size is to be taken. @@ -867,9 +913,14 @@ class Array protected: - /** \brief Method for releasing (deallocating) array data. */ + /** \brief Internal method for releasing (deallocating) array data. */ void releaseData(); + /** \brief Internal method for reallocating array elements. Used only + * from the two overloads of \ref resize. + */ + void reallocate( Index size ); + /** \brief Pointer to the data. */ Value* data = nullptr; diff --git a/src/TNL/Containers/Array.hpp b/src/TNL/Containers/Array.hpp index 6b8d1014c222671171906843dc23136bc4d8a5c6..402168d111c9ce3046eb32e9686d89054a7e9682 100644 --- a/src/TNL/Containers/Array.hpp +++ b/src/TNL/Containers/Array.hpp @@ -228,8 +228,12 @@ void Array< Value, Device, Index, Allocator >:: releaseData() { - if( this->data ) + if( this->data ) { + if( ! std::is_fundamental< ValueType >::value ) + // call the destructor of each element + Algorithms::MemoryOperations< Device >::destruct( this->data, this->size ); allocator.deallocate( this->data, this->size ); + } this->data = nullptr; this->size = 0; } @@ -240,22 +244,96 @@ template< typename Value, typename Allocator > void Array< Value, Device, Index, Allocator >:: -setSize( Index size ) +reallocate( Index size ) { TNL_ASSERT_GE( size, (Index) 0, "Array size must be non-negative." ); if( this->size == size ) return; - this->releaseData(); // Allocating zero bytes is useless. Moreover, the allocators don't behave the same way: // "operator new" returns some non-zero address, the latter returns a null pointer. - if( size > 0 ) { + if( size == 0 ) { + this->releaseData(); + return; + } + + // handle initial allocations + if( this->size == 0 ) { this->data = allocator.allocate( size ); + if( ! std::is_fundamental< ValueType >::value ) + // call the constructor of each element + Algorithms::MemoryOperations< Device >::construct( this->data, size ); + this->size = size; TNL_ASSERT_TRUE( this->data, "This should never happen - allocator did not throw on an error." ); + return; } + + // allocate an array with the correct size + Array aux( size ); + + // copy the old elements into aux + Algorithms::MemoryOperations< Device >:: + copy( aux.getData(), this->getData(), TNL::min( this->size, size ) ); + + // swap *this with aux, old data will be released + this->swap( aux ); +} + +template< typename Value, + typename Device, + typename Index, + typename Allocator > +void +Array< Value, Device, Index, Allocator >:: +resize( Index size ) +{ + // remember the old size and reallocate the array + const Index old_size = this->size; + reallocate( size ); + + if( old_size < size ) + if( ! std::is_fundamental< ValueType >::value ) + // initialize the appended elements + Algorithms::MemoryOperations< Device >::construct( this->data + old_size, size - old_size ); +} + +template< typename Value, + typename Device, + typename Index, + typename Allocator > +void +Array< Value, Device, Index, Allocator >:: +resize( Index size, const ValueType& value ) +{ + // remember the old size and reallocate the array + const Index old_size = this->size; + reallocate( size ); + + if( old_size < size ) + // copy value into the appended elements + Algorithms::MemoryOperations< Device >::construct( this->data + old_size, size - old_size, value ); +} + +template< typename Value, + typename Device, + typename Index, + typename Allocator > +void +Array< Value, Device, Index, Allocator >:: +setSize( Index size ) +{ + TNL_ASSERT_GE( size, (Index) 0, "Array size must be non-negative." ); + + if( this->size == size ) + return; + + // release data to avoid copying the elements to the new memory location + this->releaseData(); + // resize from size 0 does not copy anything, initialization is done as intended + this->resize( size ); } template< typename Value, diff --git a/src/UnitTests/Algorithms/CMakeLists.txt b/src/UnitTests/Algorithms/CMakeLists.txt index dd439fbb3f8599e681ae5f05ea13ecd72f23be08..30ea96b4da4a0a3fa41e3c031e4e14942833c781 100644 --- a/src/UnitTests/Algorithms/CMakeLists.txt +++ b/src/UnitTests/Algorithms/CMakeLists.txt @@ -1,3 +1,5 @@ +ADD_SUBDIRECTORY( Segments ) + set( COMMON_TESTS MemoryOperationsTest MultireductionTest diff --git a/src/UnitTests/Containers/Segments/CMakeLists.txt b/src/UnitTests/Algorithms/Segments/CMakeLists.txt similarity index 100% rename from src/UnitTests/Containers/Segments/CMakeLists.txt rename to src/UnitTests/Algorithms/Segments/CMakeLists.txt diff --git a/src/UnitTests/Containers/Segments/SegmentsTest.hpp b/src/UnitTests/Algorithms/Segments/SegmentsTest.hpp similarity index 100% rename from src/UnitTests/Containers/Segments/SegmentsTest.hpp rename to src/UnitTests/Algorithms/Segments/SegmentsTest.hpp diff --git a/src/UnitTests/Containers/Segments/SegmentsTest_CSR.cpp b/src/UnitTests/Algorithms/Segments/SegmentsTest_CSR.cpp similarity index 100% rename from src/UnitTests/Containers/Segments/SegmentsTest_CSR.cpp rename to src/UnitTests/Algorithms/Segments/SegmentsTest_CSR.cpp diff --git a/src/UnitTests/Containers/Segments/SegmentsTest_CSR.cu b/src/UnitTests/Algorithms/Segments/SegmentsTest_CSR.cu similarity index 100% rename from src/UnitTests/Containers/Segments/SegmentsTest_CSR.cu rename to src/UnitTests/Algorithms/Segments/SegmentsTest_CSR.cu diff --git a/src/UnitTests/Containers/Segments/SegmentsTest_CSR.h b/src/UnitTests/Algorithms/Segments/SegmentsTest_CSR.h similarity index 100% rename from src/UnitTests/Containers/Segments/SegmentsTest_CSR.h rename to src/UnitTests/Algorithms/Segments/SegmentsTest_CSR.h diff --git a/src/UnitTests/Containers/Segments/SegmentsTest_Ellpack.cpp b/src/UnitTests/Algorithms/Segments/SegmentsTest_Ellpack.cpp similarity index 100% rename from src/UnitTests/Containers/Segments/SegmentsTest_Ellpack.cpp rename to src/UnitTests/Algorithms/Segments/SegmentsTest_Ellpack.cpp diff --git a/src/UnitTests/Containers/Segments/SegmentsTest_Ellpack.cu b/src/UnitTests/Algorithms/Segments/SegmentsTest_Ellpack.cu similarity index 100% rename from src/UnitTests/Containers/Segments/SegmentsTest_Ellpack.cu rename to src/UnitTests/Algorithms/Segments/SegmentsTest_Ellpack.cu diff --git a/src/UnitTests/Containers/Segments/SegmentsTest_Ellpack.h b/src/UnitTests/Algorithms/Segments/SegmentsTest_Ellpack.h similarity index 100% rename from src/UnitTests/Containers/Segments/SegmentsTest_Ellpack.h rename to src/UnitTests/Algorithms/Segments/SegmentsTest_Ellpack.h diff --git a/src/UnitTests/Containers/Segments/SegmentsTest_SlicedEllpack.cpp b/src/UnitTests/Algorithms/Segments/SegmentsTest_SlicedEllpack.cpp similarity index 100% rename from src/UnitTests/Containers/Segments/SegmentsTest_SlicedEllpack.cpp rename to src/UnitTests/Algorithms/Segments/SegmentsTest_SlicedEllpack.cpp diff --git a/src/UnitTests/Containers/Segments/SegmentsTest_SlicedEllpack.cu b/src/UnitTests/Algorithms/Segments/SegmentsTest_SlicedEllpack.cu similarity index 100% rename from src/UnitTests/Containers/Segments/SegmentsTest_SlicedEllpack.cu rename to src/UnitTests/Algorithms/Segments/SegmentsTest_SlicedEllpack.cu diff --git a/src/UnitTests/Containers/Segments/SegmentsTest_SlicedEllpack.h b/src/UnitTests/Algorithms/Segments/SegmentsTest_SlicedEllpack.h similarity index 100% rename from src/UnitTests/Containers/Segments/SegmentsTest_SlicedEllpack.h rename to src/UnitTests/Algorithms/Segments/SegmentsTest_SlicedEllpack.h diff --git a/src/UnitTests/Containers/ArrayTest.h b/src/UnitTests/Containers/ArrayTest.h index 4b5809747a45adc992bad91458c3d6a08771000c..148d92aa659b1754cce8b45367df3d9a361d2bc6 100644 --- a/src/UnitTests/Containers/ArrayTest.h +++ b/src/UnitTests/Containers/ArrayTest.h @@ -222,6 +222,69 @@ TYPED_TEST( ArrayTest, constructorsWithAllocators ) EXPECT_EQ( a3.getElement( 1 ), 8 ); EXPECT_EQ( a3.getElement( 2 ), 9 ); EXPECT_EQ( a3.getAllocator(), allocator ); + + // test value-initialization of non-fundamental types + if( ! std::is_fundamental< typename ArrayType::ValueType >::value ) + { + const typename ArrayType::ValueType init{}; + ArrayType a( 42 ); + ASSERT_EQ( a.getSize(), 42 ); + for( int i = 0; i < a.getSize(); i++ ) + EXPECT_EQ( a.getElement( i ), init ) << "i = " << i; + } +} + +TYPED_TEST( ArrayTest, resize ) +{ + using ArrayType = typename TestFixture::ArrayType; + + ArrayType u( 42 ); + ASSERT_EQ( u.getSize(), 42 ); + for( int i = 0; i < u.getSize(); i++ ) + u.setElement( i, i ); + + // no change test + const typename ArrayType::ValueType* old_data = u.getData(); + u.resize( u.getSize() ); + EXPECT_EQ( u.getData(), old_data ); + + // shrink test + u.resize( 20 ); + ASSERT_EQ( u.getSize(), 20 ); + EXPECT_NE( u.getData(), old_data ); + for( int i = 0; i < u.getSize(); i++ ) + EXPECT_EQ( u.getElement( i ), i ); + + // expand test + const typename ArrayType::IndexType old_size = u.getSize(); + old_data = u.getData(); + u.resize( old_size * 2 ); + ASSERT_EQ( u.getSize(), old_size * 2 ); + EXPECT_NE( u.getData(), old_data ); + for( int i = 0; i < old_size; i++ ) + EXPECT_EQ( u.getElement( i ), i ); + + // expand test with initial value + const typename ArrayType::ValueType init = 3; + ArrayType v( 10 ); + v.setValue( 0 ); + v.resize( 42, init ); + ASSERT_EQ( v.getSize(), 42 ); + for( int i = 0; i < 10; i++ ) + EXPECT_EQ( v.getElement( i ), 0 ) << "i = " << i; + for( int i = 10; i < v.getSize(); i++ ) + EXPECT_EQ( v.getElement( i ), init ) << "i = " << i; + + // test value-initialization of non-fundamental types + if( ! std::is_fundamental< typename ArrayType::ValueType >::value ) + { + const typename ArrayType::ValueType init{}; + ArrayType w; + w.resize( 42 ); + ASSERT_EQ( w.getSize(), 42 ); + for( int i = 0; i < w.getSize(); i++ ) + EXPECT_EQ( w.getElement( i ), init ) << "i = " << i; + } } TYPED_TEST( ArrayTest, setSize ) @@ -373,21 +436,23 @@ TYPED_TEST( ArrayTest, setElement ) test_setElement< ArrayType >(); } -TYPED_TEST( ArrayTest, forElements ) +// test must be in a plain function because nvcc sucks (extended lambdas are +// not allowed to be defined in protected class member functions) +template< typename ArrayType > +void testArrayForEachElement() { - using ArrayType = typename TestFixture::ArrayType; using IndexType = typename ArrayType::IndexType; using ValueType = typename ArrayType::ValueType; -#if not defined HAVE_CUDA -// nvcc does not accept the following code with -// error #3068-D: The enclosing parent function ("TestBody") for an extended __host__ __device__ lambda cannot have private or protected access within its class ArrayType a( 10 ); a.forEachElement( [] __cuda_callable__ ( IndexType i, ValueType& v ) mutable { v = i; } ); for( int i = 0; i < 10; i++ ) EXPECT_EQ( a.getElement( i ), i ); -#endif +} +TYPED_TEST( ArrayTest, forElements ) +{ + testArrayForEachElement< typename TestFixture::ArrayType >(); } TYPED_TEST( ArrayTest, containsValue ) diff --git a/src/UnitTests/Containers/CMakeLists.txt b/src/UnitTests/Containers/CMakeLists.txt index efba5e50de2756a9f9335ecf18094e99689bad08..9d9e413432d33aad501c184bbf24bfae2bc41da6 100644 --- a/src/UnitTests/Containers/CMakeLists.txt +++ b/src/UnitTests/Containers/CMakeLists.txt @@ -1,5 +1,4 @@ ADD_SUBDIRECTORY( ndarray ) -ADD_SUBDIRECTORY( Segments ) set( CPP_TESTS ArrayTest diff --git a/src/UnitTests/Containers/VectorTest.h b/src/UnitTests/Containers/VectorTest.h index 136154fdcc52ea9609cf561790fa76d7730d6f90..fa99547ec325eb9a3a9b6f10c38d5ad252d6bf66 100644 --- a/src/UnitTests/Containers/VectorTest.h +++ b/src/UnitTests/Containers/VectorTest.h @@ -80,15 +80,14 @@ TYPED_TEST( VectorTest, constructors ) } -TYPED_TEST( VectorTest, reduceElements ) +// test must be in a plain function because nvcc sucks (extended lambdas are +// not allowed to be defined in protected class member functions) +template< typename VectorType > +void testVectorReduceElements() { - using VectorType = typename TestFixture::VectorType; using IndexType = typename VectorType::IndexType; using ValueType = typename VectorType::ValueType; -#if not defined HAVE_CUDA -// nvcc does not accept the following code with -// error #3068-D: The enclosing parent function ("TestBody") for an extended __host__ __device__ lambda cannot have private or protected access within its class VectorType a( 10 ); a.forEachElement( [=] __cuda_callable__ ( IndexType i, ValueType& v ) mutable { v = 1; } ); auto fetch = [] __cuda_callable__ ( IndexType i, ValueType& v ) -> ValueType { return v; }; @@ -100,7 +99,10 @@ TYPED_TEST( VectorTest, reduceElements ) auto const_fetch = [] __cuda_callable__ ( IndexType i, const ValueType& v ) -> ValueType { return v; }; EXPECT_EQ( b.reduceEachElement( const_fetch, reduce, ( ValueType ) 0.0 ), b.getSize() ); -#endif +} +TYPED_TEST( VectorTest, reduceElements ) +{ + testVectorReduceElements< typename TestFixture::VectorType >(); } TEST( VectorSpecialCasesTest, defaultConstructors )