diff --git a/CMakeLists.txt b/CMakeLists.txt index 869aeed0..c8c32793 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -5,7 +5,7 @@ # SPDX-License-Identifier: (BSD-3-Clause) ################################################################################################### -cmake_minimum_required( VERSION 3.9 ) +cmake_minimum_required( VERSION 3.24 ) # Set version number set( LVARRAY_VERSION_MAJOR 0 ) @@ -77,11 +77,11 @@ blt_list_append( TO lvarray_dependencies ELEMENTS chai IF ENABLE_CHAI ) blt_list_append( TO lvarray_dependencies ELEMENTS RAJA ) -blt_list_append( TO lvarray_dependencies ELEMENTS umpire IF ENABLE_UMPIRE ) +blt_list_append( TO lvarray_dependencies ELEMENTS umpire::umpire IF ENABLE_UMPIRE ) # Ignore umpire warnings as errors by using -isystem flag if (ENABLE_UMPIRE) - blt_convert_to_system_includes(TARGET umpire) + blt_convert_to_system_includes(TARGET umpire::umpire) endif() blt_list_append( TO lvarray_dependencies ELEMENTS cuda IF ENABLE_CUDA ) @@ -116,5 +116,3 @@ endif() if( ENABLE_DOCS ) add_subdirectory( docs ) endif() - - diff --git a/cmake/SetupTPL.cmake b/cmake/SetupTPL.cmake index bff94834..b0bdb261 100644 --- a/cmake/SetupTPL.cmake +++ b/cmake/SetupTPL.cmake @@ -44,7 +44,7 @@ if(ENABLE_UMPIRE) find_package(umpire REQUIRED PATHS ${UMPIRE_DIR}) - set(thirdPartyLibs ${thirdPartyLibs} umpire) + set(thirdPartyLibs ${thirdPartyLibs} umpire::umpire) else() message(STATUS "Not using Umpire.") endif() diff --git a/src/input.hpp b/src/input.hpp index 12cc033a..32146c0f 100644 --- a/src/input.hpp +++ b/src/input.hpp @@ -294,15 +294,17 @@ static void stringToArray( Array< T, NDIM, PERMUTATION, INDEX_TYPE, BUFFER_TYPE // we also need to add a ' ' in front of any '}' otherwise the // stringstream::operator>> will grab the } - for( std::string::size_type a=0; a::Read( array.toSlice(), array.dims(), strstream ); } diff --git a/unitTests/testArray1DOfArray1D.cpp b/unitTests/testArray1DOfArray1D.cpp index 51823a3f..4474396e 100644 --- a/unitTests/testArray1DOfArray1D.cpp +++ b/unitTests/testArray1DOfArray1D.cpp @@ -389,7 +389,7 @@ class Array1DOfArray1DOfArrayView2DTest : public ::testing::Test { for( IndexType c = 0; c < nestedView[ i ][ j ].size( 1 ); ++c ) { - PORTABLE_EXPECT_EQ( nestedView[ i ][ j ]( r, c ), deviceTouchedValue( i, j, r, c ) ); + PORTABLE_EXPECT_NEAR( nestedView[ i ][ j ]( r, c ), deviceTouchedValue( i, j, r, c ), T( 1e-12 ) ); } } } @@ -423,7 +423,7 @@ class Array1DOfArray1DOfArrayView2DTest : public ::testing::Test { for( IndexType c = 0; c < nestedView[ i ][ j ].size( 1 ); ++c ) { - PORTABLE_EXPECT_EQ( nestedView[ i ][ j ]( r, c ), initialValue( i, j, r, c ) ); + PORTABLE_EXPECT_NEAR( nestedView[ i ][ j ]( r, c ), initialValue( i, j, r, c ), T( 1e-12 ) ); } } } @@ -440,7 +440,7 @@ class Array1DOfArray1DOfArrayView2DTest : public ::testing::Test { for( IndexType c = 0; c < nestedView[ i ][ j ].size( 1 ); ++c ) { - PORTABLE_EXPECT_EQ( nestedView[ i ][ j ]( r, c ), hostTouchedValue( i, j, r, c ) ); + PORTABLE_EXPECT_NEAR( nestedView[ i ][ j ]( r, c ), hostTouchedValue( i, j, r, c ), T( 1e-12 ) ); } } } diff --git a/unitTests/testArrayView.hpp b/unitTests/testArrayView.hpp index 961723b8..c71d7f75 100644 --- a/unitTests/testArrayView.hpp +++ b/unitTests/testArrayView.hpp @@ -435,9 +435,9 @@ class ArrayViewPolicyTest : public ArrayViewTest< typename ARRAY_POLICY_PAIR::fi array->template setValues< POLICY >( value ); ViewTypeConst const view = array->toViewConst(); - forall< POLICY >( array->size(), [view, value] LVARRAY_HOST_DEVICE ( INDEX_TYPE const i ) + forall< POLICY >( array->size(), [view] LVARRAY_HOST_DEVICE ( INDEX_TYPE const i ) { - PORTABLE_EXPECT_EQ( view.data()[ i ], value ); + PORTABLE_EXPECT_EQ( view.data()[ i ], T( 3.14 ) ); } ); EXPECT_EQ( array->size(), totalSize ); diff --git a/unitTests/testMath.cpp b/unitTests/testMath.cpp index cd04b250..a0b88cb7 100644 --- a/unitTests/testMath.cpp +++ b/unitTests/testMath.cpp @@ -373,7 +373,7 @@ struct TestMath2 : public ::testing::Test } }; -#if defined( LVARRAY_USE_CUDA ) || defined( LVARRAY_USE_HIP ) +#if defined( LVARRAY_USE_CUDA ) using TestMath2Types = ::testing::Types< std::pair< __half2, parallelDevicePolicy< 32 > > diff --git a/unitTests/testMemcpy.cpp b/unitTests/testMemcpy.cpp index 0e44243d..c82d569a 100644 --- a/unitTests/testMemcpy.cpp +++ b/unitTests/testMemcpy.cpp @@ -65,7 +65,7 @@ void testAsyncMemcpy1D() Array< int, 1, RAJA::PERM_I, std::ptrdiff_t, BUFFER_TYPE > y( x.size() ); camp::resources::Event e = memcpy( host, y.toSlice(), x.toSliceConst() ); - host.wait_for( &e ); + host.wait_for( e ); for( std::ptrdiff_t i = 0; i < x.size(); ++i ) { @@ -78,7 +78,7 @@ void testAsyncMemcpy1D() } e = memcpy< 0, 0 >( host, y, {}, x.toViewConst(), {} ); - host.wait_for( &e ); + host.wait_for( e ); for( std::ptrdiff_t i = 0; i < x.size(); ++i ) { @@ -162,7 +162,7 @@ void testMemcpyDevice() forall< RAJA::cuda_exec< 32 > >( y.size(), [yPtr] LVARRAY_DEVICE ( std::ptrdiff_t const i ) { - PORTABLE_EXPECT_EQ( yPtr[ i ], i ); + PORTABLE_DEVICE_EXPECT_EQ( yPtr[ i ], i ); yPtr[ i ] *= 2; } ); @@ -194,7 +194,7 @@ void testMemcpyDevice() template< template< typename > class BUFFER_TYPE > void testAsyncMemcpyDevice() { - camp::resources::Resource stream{ camp::resources::Cuda{} }; + camp::resources::Resource stream{ camp::resources::Cuda::get_default() }; Array< int, 1, RAJA::PERM_I, std::ptrdiff_t, BUFFER_TYPE > x( 100 ); @@ -208,16 +208,18 @@ void testAsyncMemcpyDevice() int * yPtr = y.data(); camp::resources::Event e = memcpy< 0, 0 >( stream, y.toView(), {}, x.toViewConst(), {} ); - stream.wait_for( &e ); + stream.wait_for( e ); - forall< RAJA::cuda_exec< 32 > >( y.size(), [yPtr] LVARRAY_DEVICE ( std::ptrdiff_t const i ) + RAJA::forall< RAJA::cuda_exec< 32 > >( stream.get< camp::resources::Cuda >(), + RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, y.size() ), + [yPtr] LVARRAY_DEVICE ( std::ptrdiff_t const i ) { - PORTABLE_EXPECT_EQ( yPtr[ i ], i ); + PORTABLE_DEVICE_EXPECT_EQ( yPtr[ i ], i ); yPtr[ i ] *= 2; } ); e = memcpy< 0, 0 >( stream, x, {}, y.toViewConst(), {} ); - stream.wait_for( &e ); + stream.wait_for( e ); for( std::ptrdiff_t i = 0; i < x.size(); ++i ) { @@ -229,13 +231,15 @@ void testAsyncMemcpyDevice() y.move( MemorySpace::host ); ArrayView< int, 1, 0, std::ptrdiff_t, BUFFER_TYPE > const yView = y.toView(); - forall< RAJA::cuda_exec< 32 > >( y.size(), [yView] LVARRAY_DEVICE ( std::ptrdiff_t const i ) + RAJA::forall< RAJA::cuda_exec< 32 > >( stream.get< camp::resources::Cuda >(), + RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, y.size() ), + [yView] LVARRAY_DEVICE ( std::ptrdiff_t const i ) { yView[ i ] = -i; } ); e = memcpy< 0, 0 >( stream, x, {}, y.toViewConst(), {} ); - stream.wait_for( &e ); + stream.wait_for( e ); for( std::ptrdiff_t i = 0; i < x.size(); ++i ) { @@ -262,7 +266,7 @@ void testMemcpyDevice() forall< RAJA::hip_exec< 32 > >( y.size(), [yPtr] LVARRAY_DEVICE ( std::ptrdiff_t const i ) { - PORTABLE_EXPECT_EQ( yPtr[ i ], i ); + PORTABLE_DEVICE_EXPECT_EQ( yPtr[ i ], i ); yPtr[ i ] *= 2; } ); @@ -294,7 +298,7 @@ void testMemcpyDevice() template< template< typename > class BUFFER_TYPE > void testAsyncMemcpyDevice() { - camp::resources::Resource stream{ camp::resources::Hip{} }; + camp::resources::Resource stream{ camp::resources::Hip::get_default() }; Array< int, 1, RAJA::PERM_I, std::ptrdiff_t, BUFFER_TYPE > x( 100 ); @@ -308,16 +312,18 @@ void testAsyncMemcpyDevice() int * yPtr = y.data(); camp::resources::Event e = memcpy< 0, 0 >( stream, y.toView(), {}, x.toViewConst(), {} ); - stream.wait_for( &e ); + stream.wait_for( e ); - forall< RAJA::hip_exec< 32 > >( y.size(), [yPtr] LVARRAY_DEVICE ( std::ptrdiff_t const i ) + RAJA::forall< RAJA::hip_exec< 32 > >( stream.get< camp::resources::Hip >(), + RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, y.size() ), + [yPtr] LVARRAY_DEVICE ( std::ptrdiff_t const i ) { - PORTABLE_EXPECT_EQ( yPtr[ i ], i ); + PORTABLE_DEVICE_EXPECT_EQ( yPtr[ i ], i ); yPtr[ i ] *= 2; } ); e = memcpy< 0, 0 >( stream, x, {}, y.toViewConst(), {} ); - stream.wait_for( &e ); + stream.wait_for( e ); for( std::ptrdiff_t i = 0; i < x.size(); ++i ) { @@ -329,13 +335,15 @@ void testAsyncMemcpyDevice() y.move( MemorySpace::host ); ArrayView< int, 1, 0, std::ptrdiff_t, BUFFER_TYPE > const yView = y.toView(); - forall< RAJA::hip_exec< 32 > >( y.size(), [yView] LVARRAY_DEVICE ( std::ptrdiff_t const i ) + RAJA::forall< RAJA::hip_exec< 32 > >( stream.get< camp::resources::Hip >(), + RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, y.size() ), + [yView] LVARRAY_DEVICE ( std::ptrdiff_t const i ) { yView[ i ] = -i; } ); e = memcpy< 0, 0 >( stream, x, {}, y.toViewConst(), {} ); - stream.wait_for( &e ); + stream.wait_for( e ); for( std::ptrdiff_t i = 0; i < x.size(); ++i ) { diff --git a/unitTests/testStackArray.cpp b/unitTests/testStackArray.cpp index e29206ab..224f8a09 100644 --- a/unitTests/testStackArray.cpp +++ b/unitTests/testStackArray.cpp @@ -175,36 +175,43 @@ class StackArrayCaptureTest : public StackArrayTest< typename PERMUTATION_POLICY /// This needs to use the parallelDevice policy because you can't nest host-device lambdas. static void resizeMultipleInLambda() { - INDEX_TYPE dims[ NDIM ]; - for( int i = 0; i < NDIM; ++i ) - { dims[ i ] = 8; } + if constexpr ( std::is_same< POLICY, serialPolicy >::value ) + { + GTEST_SKIP() << "Nested device lambdas are only supported by a device execution policy."; + return; + } + else + { + INDEX_TYPE dims[ NDIM ]; + for( int i = 0; i < NDIM; ++i ) + { dims[ i ] = 8; } - INDEX_TYPE const capacity = CAPACITY; - forall< POLICY >( 10, [dims, capacity] LVARRAY_DEVICE ( int ) - { - StackArray< int, NDIM, PERMUTATION, INDEX_TYPE, CAPACITY > array; - PORTABLE_EXPECT_EQ( array.size(), 0 ); - PORTABLE_EXPECT_EQ( array.capacity(), capacity ); + forall< POLICY >( 10, [dims] LVARRAY_DEVICE ( int ) + { + StackArray< int, NDIM, PERMUTATION, INDEX_TYPE, CAPACITY > array; + PORTABLE_DEVICE_EXPECT_EQ( array.size(), 0 ); + PORTABLE_DEVICE_EXPECT_EQ( array.capacity(), CAPACITY ); - array.resize( NDIM, dims ); + array.resize( NDIM, dims ); - for( int i = 0; i < NDIM; ++i ) - { PORTABLE_EXPECT_EQ( array.size( i ), 8 ); } + for( int i = 0; i < NDIM; ++i ) + { PORTABLE_DEVICE_EXPECT_EQ( array.size( i ), 8 ); } - PORTABLE_EXPECT_EQ( array.size(), capacity ); + PORTABLE_DEVICE_EXPECT_EQ( array.size(), CAPACITY ); - forValuesInSliceWithIndices( array.toSlice(), SetValue() ); + forValuesInSliceWithIndices( array.toSlice(), SetValue() ); - array.resize( 2 ); + array.resize( 2 ); - PORTABLE_EXPECT_EQ( array.size( 0 ), 2 ); - for( int i = 1; i < NDIM; ++i ) - { PORTABLE_EXPECT_EQ( array.size( i ), 8 ); } + PORTABLE_DEVICE_EXPECT_EQ( array.size( 0 ), 2 ); + for( int i = 1; i < NDIM; ++i ) + { PORTABLE_DEVICE_EXPECT_EQ( array.size( i ), 8 ); } - PORTABLE_EXPECT_EQ( array.size(), array.capacity() / 4 ); + PORTABLE_DEVICE_EXPECT_EQ( array.size(), array.capacity() / 4 ); - forValuesInSliceWithIndices( array.toSlice(), CheckValue() ); - } ); + forValuesInSliceWithIndices( array.toSlice(), CheckValue() ); + } ); + } } template< typename _PERMUTATION=PERMUTATION > @@ -212,7 +219,7 @@ class StackArrayCaptureTest : public StackArrayTest< typename PERMUTATION_POLICY sizedConstructorInLambda() { INDEX_TYPE const capacity = CAPACITY; - forall< POLICY >( 10, [capacity] LVARRAY_DEVICE ( int ) + forall< POLICY >( 10, [capacity] LVARRAY_HOST_DEVICE ( int ) { StackArray< int, NDIM, PERMUTATION, INDEX_TYPE, CAPACITY > array( CAPACITY ); PORTABLE_EXPECT_EQ( array.capacity(), capacity ); @@ -227,7 +234,7 @@ class StackArrayCaptureTest : public StackArrayTest< typename PERMUTATION_POLICY { INDEX_TYPE const capacity = CAPACITY; int const size = 8; - forall< POLICY >( 10, [capacity, size] LVARRAY_DEVICE ( int ) + forall< POLICY >( 10, [capacity, size] LVARRAY_HOST_DEVICE ( int ) { StackArray< int, NDIM, PERMUTATION, INDEX_TYPE, CAPACITY > array( size - 1, size ); PORTABLE_EXPECT_EQ( array.capacity(), capacity ); @@ -243,7 +250,7 @@ class StackArrayCaptureTest : public StackArrayTest< typename PERMUTATION_POLICY { INDEX_TYPE const capacity = CAPACITY; int const size = 8; - forall< POLICY >( 10, [capacity, size] LVARRAY_DEVICE ( int ) + forall< POLICY >( 10, [capacity, size] LVARRAY_HOST_DEVICE ( int ) { StackArray< int, NDIM, PERMUTATION, INDEX_TYPE, CAPACITY > array( size - 2, size - 1, size ); PORTABLE_EXPECT_EQ( array.capacity(), capacity ); diff --git a/unitTests/testTensorOpsCommon.hpp b/unitTests/testTensorOpsCommon.hpp index 39921996..1feea2d9 100644 --- a/unitTests/testTensorOpsCommon.hpp +++ b/unitTests/testTensorOpsCommon.hpp @@ -119,5 +119,22 @@ randomValue( T const maxVal, std::mt19937_64 & gen ) } \ } while ( false ) +#define CHECK_NEAR_2D( N, M, A, RESULT, EPSILON ) \ + tensorOps::internal::checkSizes< N, M >( A ); \ + tensorOps::internal::checkSizes< N, M >( RESULT ); \ + do \ + { \ + for( std::ptrdiff_t _i = 0; _i < N; ++_i ) \ + { \ + for( std::ptrdiff_t _j = 0; _j < M; ++_j ) \ + { \ + if( std::is_integral< std::remove_reference_t< decltype( A[ _i ][ _j ] ) > >::value ) \ + { PORTABLE_EXPECT_EQ( A[ _i ][ _j ], RESULT[ _i ][ _j ] ); } \ + else \ + { PORTABLE_EXPECT_NEAR( A[ _i ][ _j ], RESULT[ _i ][ _j ], EPSILON ); } \ + } \ + } \ + } while ( false ) + } // namespace testing } // namespace LvArray diff --git a/unitTests/testTensorOpsEigen.cpp b/unitTests/testTensorOpsEigen.cpp index 2c556ec7..973b0698 100644 --- a/unitTests/testTensorOpsEigen.cpp +++ b/unitTests/testTensorOpsEigen.cpp @@ -104,7 +104,7 @@ class TestEigendecomposition : public ::testing::Test ArrayViewT< FLOAT const, 2, 1 > const eigenvalues = m_eigenvalues.toViewConst(); ArrayViewT< FLOAT, 2, 1 > const expectedEigenvalues = m_expectedEigenvalues.toView(); ArrayViewT< FLOAT const, 3, 2 > const eigenvectors = m_eigenvectors.toViewConst(); - forall< serialPolicy >( matrices.size( 0 ), [=, &relativeDiffs] ( std::ptrdiff_t const i ) + forall< serialPolicy >( matrices.size( 0 ), [=, &relativeDiffs, this] ( std::ptrdiff_t const i ) { if( iteration > 0 ) { @@ -174,7 +174,7 @@ class TestEigendecomposition : public ::testing::Test ArrayViewT< FLOAT, 2, 1 > const eigenvalues = m_eigenvalues.toView(); ArrayViewT< FLOAT, 2, 1 > const expectedEigenvalues = m_expectedEigenvalues.toView(); ArrayViewT< FLOAT, 3, 2 > const eigenvectors = m_eigenvectors.toView(); - forall< serialPolicy >( matrices.size( 0 ), [=] ( std::ptrdiff_t const i ) + forall< serialPolicy >( matrices.size( 0 ), [=, this] ( std::ptrdiff_t const i ) { // Since we're constructing the matrix we know the eigenvalues beforehand. for( INDEX_TYPE j = 0; j < M; ++j ) diff --git a/unitTests/testTensorOpsOneSize.cpp b/unitTests/testTensorOpsOneSize.cpp index 78946638..efdbadf5 100644 --- a/unitTests/testTensorOpsOneSize.cpp +++ b/unitTests/testTensorOpsOneSize.cpp @@ -83,7 +83,7 @@ class OneSizeTest : public ::testing::Test T vectorA_local[ N ]; fill( vectorA_local, aSeed ); tensorOps::scale< N >( vectorA_local, scale ); - CHECK_EQUALITY_1D( N, vectorA_local, result ); + CHECK_NEAR_1D( N, vectorA_local, result, 100 * NumericLimits< T >::epsilon ); } ); } @@ -264,7 +264,12 @@ class OneSizeTest : public ::testing::Test #define _TEST( a, b ) \ tensorOps::scaledCopy< N >( a, b, scale ); \ for( std::ptrdiff_t i = 0; i < N; ++i ) \ - { PORTABLE_EXPECT_EQ( a[ i ], scale * b[ i ] ); } \ + { \ + if constexpr ( std::is_integral< T >::value ) \ + { PORTABLE_EXPECT_EQ( a[ i ], scale * b[ i ] ); } \ + else \ + { PORTABLE_EXPECT_NEAR( a[ i ], scale * b[ i ], 100 * NumericLimits< T >::epsilon ); } \ + } \ fill( a, aSeed ) #define _TEST_PERMS( a, b0, b1, b2 ) \ diff --git a/unitTests/testTensorOpsTwoSizes.hpp b/unitTests/testTensorOpsTwoSizes.hpp index 5492b2b5..d5da5e4f 100644 --- a/unitTests/testTensorOpsTwoSizes.hpp +++ b/unitTests/testTensorOpsTwoSizes.hpp @@ -82,18 +82,18 @@ class TwoSizesTest : public ::testing::Test forall< POLICY >( 1, [scale, result, matrixA_IJK, matrixA_IKJ, matrixA_KJI, aSeed] LVARRAY_HOST_DEVICE ( int ) { tensorOps::scale< N, M >( matrixA_IJK[ 0 ], scale ); - CHECK_EQUALITY_2D( N, M, matrixA_IJK[ 0 ], result ); + CHECK_NEAR_2D( N, M, matrixA_IJK[ 0 ], result, 100 * NumericLimits< T >::epsilon ); tensorOps::scale< N, M >( matrixA_IKJ[ 0 ], scale ); - CHECK_EQUALITY_2D( N, M, matrixA_IKJ[ 0 ], result ); + CHECK_NEAR_2D( N, M, matrixA_IKJ[ 0 ], result, 100 * NumericLimits< T >::epsilon ); tensorOps::scale< N, M >( matrixA_KJI[ 0 ], scale ); - CHECK_EQUALITY_2D( N, M, matrixA_KJI[ 0 ], result ); + CHECK_NEAR_2D( N, M, matrixA_KJI[ 0 ], result, 100 * NumericLimits< T >::epsilon ); T matrix_local[ N ][ M ]; fill( matrix_local, aSeed ); tensorOps::scale< N, M >( matrix_local, scale ); - CHECK_EQUALITY_2D( N, M, matrix_local, result ); + CHECK_NEAR_2D( N, M, matrix_local, result, 100 * NumericLimits< T >::epsilon ); } ); } @@ -585,7 +585,7 @@ class TwoSizesTest : public ::testing::Test #define _TEST( dstMatrix, srcMatrix ) \ fill( dstMatrix, matrixSeed ); \ tensorOps::scaledCopy< N, M >( dstMatrix, srcMatrix, scale ); \ - CHECK_EQUALITY_2D( N, M, dstMatrix, result ) + CHECK_NEAR_2D( N, M, dstMatrix, result, 100 * NumericLimits< T >::epsilon ) #define _TEST_PERMS( dstMatrix, srcMatrix0, srcMatrix1, srcMatrix2, srcMatrix3 ) \ _TEST( dstMatrix, srcMatrix0 ); \ @@ -689,7 +689,7 @@ class TwoSizesTest : public ::testing::Test #define _TEST( dstMatrix, srcMatrix ) \ fill( dstMatrix, matrixSeed ); \ tensorOps::scaledAdd< N, M >( dstMatrix, srcMatrix, scale ); \ - CHECK_EQUALITY_2D( N, M, dstMatrix, result ); \ + CHECK_NEAR_2D( N, M, dstMatrix, result, 100 * NumericLimits< T >::epsilon ); \ #define _TEST_PERMS( dstMatrix, srcMatrix0, srcMatrix1, srcMatrix2, srcMatrix3 ) \ _TEST( dstMatrix, srcMatrix0 ); \ diff --git a/unitTests/testTensorOpsTwoSizes1.cpp b/unitTests/testTensorOpsTwoSizes1.cpp index 96ac793c..72e7edf2 100644 --- a/unitTests/testTensorOpsTwoSizes1.cpp +++ b/unitTests/testTensorOpsTwoSizes1.cpp @@ -83,18 +83,18 @@ class TwoSizesTest : public ::testing::Test forall< POLICY >( 1, [scale, result, matrixA_IJK, matrixA_IKJ, matrixA_KJI, aSeed] LVARRAY_HOST_DEVICE ( int ) { tensorOps::scale< N, M >( matrixA_IJK[ 0 ], scale ); - CHECK_EQUALITY_2D( N, M, matrixA_IJK[ 0 ], result ); + CHECK_NEAR_2D( N, M, matrixA_IJK[ 0 ], result, 100 * NumericLimits< T >::epsilon ); tensorOps::scale< N, M >( matrixA_IKJ[ 0 ], scale ); - CHECK_EQUALITY_2D( N, M, matrixA_IKJ[ 0 ], result ); + CHECK_NEAR_2D( N, M, matrixA_IKJ[ 0 ], result, 100 * NumericLimits< T >::epsilon ); tensorOps::scale< N, M >( matrixA_KJI[ 0 ], scale ); - CHECK_EQUALITY_2D( N, M, matrixA_KJI[ 0 ], result ); + CHECK_NEAR_2D( N, M, matrixA_KJI[ 0 ], result, 100 * NumericLimits< T >::epsilon ); T matrix_local[ N ][ M ]; fill( matrix_local, aSeed ); tensorOps::scale< N, M >( matrix_local, scale ); - CHECK_EQUALITY_2D( N, M, matrix_local, result ); + CHECK_NEAR_2D( N, M, matrix_local, result, 100 * NumericLimits< T >::epsilon ); } ); } @@ -586,7 +586,7 @@ class TwoSizesTest : public ::testing::Test #define _TEST( dstMatrix, srcMatrix ) \ fill( dstMatrix, matrixSeed ); \ tensorOps::scaledCopy< N, M >( dstMatrix, srcMatrix, scale ); \ - CHECK_EQUALITY_2D( N, M, dstMatrix, result ) + CHECK_NEAR_2D( N, M, dstMatrix, result, 100 * NumericLimits< T >::epsilon ) #define _TEST_PERMS( dstMatrix, srcMatrix0, srcMatrix1, srcMatrix2, srcMatrix3 ) \ _TEST( dstMatrix, srcMatrix0 ); \ @@ -690,7 +690,7 @@ class TwoSizesTest : public ::testing::Test #define _TEST( dstMatrix, srcMatrix ) \ fill( dstMatrix, matrixSeed ); \ tensorOps::scaledAdd< N, M >( dstMatrix, srcMatrix, scale ); \ - CHECK_EQUALITY_2D( N, M, dstMatrix, result ); \ + CHECK_NEAR_2D( N, M, dstMatrix, result, 100 * NumericLimits< T >::epsilon ); \ #define _TEST_PERMS( dstMatrix, srcMatrix0, srcMatrix1, srcMatrix2, srcMatrix3 ) \ _TEST( dstMatrix, srcMatrix0 ); \ diff --git a/unitTests/testUtils.hpp b/unitTests/testUtils.hpp index b5486988..d26643ba 100644 --- a/unitTests/testUtils.hpp +++ b/unitTests/testUtils.hpp @@ -10,6 +10,8 @@ // Source includes #include "LvArrayConfig.hpp" #include "Macros.hpp" +#include "limits.hpp" +#include "math.hpp" #include "MallocBuffer.hpp" @@ -118,7 +120,6 @@ LAYOUT const & getRAJAViewLayout( RAJA::View< T, LAYOUT > const & view ) #endif } - #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) #define PORTABLE_EXPECT_EQ( L, R ) LVARRAY_ERROR_IF_NE( L, R ) #define PORTABLE_EXPECT_NEAR( L, R, EPSILON ) LVARRAY_ERROR_IF_GE_MSG( math::abs( ( L ) -( R ) ), EPSILON, \ @@ -129,6 +130,18 @@ LAYOUT const & getRAJAViewLayout( RAJA::View< T, LAYOUT > const & view ) STRINGIZE( L ) " = " << ( L ) << "\n" << STRINGIZE( R ) " = " << ( R ); #endif +// A device-only lambda is compiled by HIP in a host pass as well. GoogleTest +// assertions are host-only and cannot be used in that lambda, so use a plain +// assertion for checks that are only made on the device. +#define PORTABLE_DEVICE_EXPECT_EQ( L, R ) \ + do \ + { \ + if( !( ( L ) == ( R ) ) ) \ + { \ + assert( false && "Device assertion failed" ); \ + } \ + } while( false ) + // Comparator that compares a std::pair by it's first object. template< class A, class B, class COMP=std::less< A > > struct PairComp