From 4b2bbb818605cf82a04947ec5607bd93cb2e51c0 Mon Sep 17 00:00:00 2001 From: "Victor A. P. Magri" Date: Fri, 21 Aug 2026 06:12:51 -0400 Subject: [PATCH 1/7] Fix C++20 test builds --- unitTests/testMemcpy.cpp | 4 ++-- unitTests/testTensorOpsEigen.cpp | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/unitTests/testMemcpy.cpp b/unitTests/testMemcpy.cpp index 0e44243d..313cb173 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 ) { 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 ) From 4c3df7f7bb341a974ce98b66ca5118454b779fdf Mon Sep 17 00:00:00 2001 From: "Victor A. P. Magri" Date: Fri, 21 Aug 2026 06:16:51 -0400 Subject: [PATCH 2/7] Fix device test builds --- unitTests/testMemcpy.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/unitTests/testMemcpy.cpp b/unitTests/testMemcpy.cpp index 313cb173..3f045461 100644 --- a/unitTests/testMemcpy.cpp +++ b/unitTests/testMemcpy.cpp @@ -208,7 +208,7 @@ 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 ) { @@ -217,7 +217,7 @@ void testAsyncMemcpyDevice() } ); 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 ) { @@ -235,7 +235,7 @@ void testAsyncMemcpyDevice() } ); 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 ) { @@ -308,7 +308,7 @@ 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 ) { @@ -317,7 +317,7 @@ void testAsyncMemcpyDevice() } ); 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 ) { @@ -335,7 +335,7 @@ void testAsyncMemcpyDevice() } ); 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 ) { From 217232ee690fed5809b832a3a35407f7f94a7bbe Mon Sep 17 00:00:00 2001 From: "Victor A. P. Magri" Date: Fri, 21 Aug 2026 06:39:52 -0400 Subject: [PATCH 3/7] Avoid parser string overlap warning --- src/input.hpp | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) 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 ); } From 9084738a6dde1fd3cbc127e5cb7d11406b30f6c9 Mon Sep 17 00:00:00 2001 From: "Victor A. P. Magri" Date: Fri, 21 Aug 2026 08:47:19 -0400 Subject: [PATCH 4/7] Fix portable device tests --- unitTests/testArray1DOfArray1D.cpp | 6 ++--- unitTests/testArrayView.hpp | 4 +-- unitTests/testMath.cpp | 2 +- unitTests/testMemcpy.cpp | 28 +++++++++++++-------- unitTests/testStackArray.cpp | 37 +++++++++++++++++----------- unitTests/testTensorOpsCommon.hpp | 17 +++++++++++++ unitTests/testTensorOpsOneSize.cpp | 9 +++++-- unitTests/testTensorOpsTwoSizes.hpp | 12 ++++----- unitTests/testTensorOpsTwoSizes1.cpp | 12 ++++----- unitTests/testUtils.hpp | 33 ++++++++++++++++++++++++- 10 files changed, 114 insertions(+), 46 deletions(-) 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 3f045461..5ebdcdf3 100644 --- a/unitTests/testMemcpy.cpp +++ b/unitTests/testMemcpy.cpp @@ -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 ); @@ -210,9 +210,11 @@ void testAsyncMemcpyDevice() camp::resources::Event e = memcpy< 0, 0 >( stream, y.toView(), {}, x.toViewConst(), {} ); 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; } ); @@ -229,7 +231,9 @@ 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; } ); @@ -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 ); @@ -310,9 +314,11 @@ void testAsyncMemcpyDevice() camp::resources::Event e = memcpy< 0, 0 >( stream, y.toView(), {}, x.toViewConst(), {} ); 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; } ); @@ -329,7 +335,9 @@ 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; } ); diff --git a/unitTests/testStackArray.cpp b/unitTests/testStackArray.cpp index e29206ab..7c6b6ad7 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 ) + forall< POLICY >( 10, [dims] LVARRAY_DEVICE ( int ) { StackArray< int, NDIM, PERMUTATION, INDEX_TYPE, CAPACITY > array; - PORTABLE_EXPECT_EQ( array.size(), 0 ); - PORTABLE_EXPECT_EQ( array.capacity(), capacity ); + PORTABLE_DEVICE_EXPECT_EQ( array.size(), 0 ); + PORTABLE_DEVICE_EXPECT_EQ( array.capacity(), CAPACITY ); array.resize( NDIM, dims ); for( int i = 0; i < NDIM; ++i ) - { PORTABLE_EXPECT_EQ( array.size( i ), 8 ); } + { 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() ); array.resize( 2 ); - PORTABLE_EXPECT_EQ( array.size( 0 ), 2 ); + PORTABLE_DEVICE_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( i ), 8 ); } - PORTABLE_EXPECT_EQ( array.size(), array.capacity() / 4 ); + PORTABLE_DEVICE_EXPECT_EQ( array.size(), array.capacity() / 4 ); 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/testTensorOpsOneSize.cpp b/unitTests/testTensorOpsOneSize.cpp index 78946638..1fe2a5a4 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..b110aa47 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,9 +120,26 @@ LAYOUT const & getRAJAViewLayout( RAJA::View< T, LAYOUT > const & view ) #endif } +template< typename LHS, typename RHS > +LVARRAY_HOST_DEVICE inline bool portableEqual( LHS const & lhs, RHS const & rhs ) +{ + using LhsValue = std::decay_t< LHS >; + using RhsValue = std::decay_t< RHS >; + if constexpr( std::is_floating_point< LhsValue >::value || std::is_floating_point< RhsValue >::value ) + { + using FloatingPoint = std::common_type_t< LhsValue, RhsValue >; + return math::abs( static_cast< FloatingPoint >( lhs ) - static_cast< FloatingPoint >( rhs ) ) <= NumericLimits< FloatingPoint >::epsilon; + } + else + { + return lhs == rhs; + } +} + #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) -#define PORTABLE_EXPECT_EQ( L, R ) LVARRAY_ERROR_IF_NE( L, R ) +#define PORTABLE_EXPECT_EQ( L, R ) LVARRAY_ERROR_IF( !portableEqual( ( L ), ( R ) ), \ + "Expected " STRINGIZE( L ) " == " STRINGIZE( R ) ) #define PORTABLE_EXPECT_NEAR( L, R, EPSILON ) LVARRAY_ERROR_IF_GE_MSG( math::abs( ( L ) -( R ) ), EPSILON, \ STRINGIZE( L ) " = " << ( L ) << "\n" << STRINGIZE( R ) " = " << ( R ) ); #else @@ -129,6 +148,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 From e05731d1897b8c668e15e426ed9e1fb10d4d5638 Mon Sep 17 00:00:00 2001 From: "Victor A. P. Magri" Date: Fri, 21 Aug 2026 11:10:06 -0400 Subject: [PATCH 5/7] Format device tests --- unitTests/testMemcpy.cpp | 16 +++++++------- unitTests/testStackArray.cpp | 34 +++++++++++++++--------------- unitTests/testTensorOpsOneSize.cpp | 2 +- unitTests/testUtils.hpp | 2 +- 4 files changed, 27 insertions(+), 27 deletions(-) diff --git a/unitTests/testMemcpy.cpp b/unitTests/testMemcpy.cpp index 5ebdcdf3..c82d569a 100644 --- a/unitTests/testMemcpy.cpp +++ b/unitTests/testMemcpy.cpp @@ -211,8 +211,8 @@ void testAsyncMemcpyDevice() stream.wait_for( e ); 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 ) + RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, y.size() ), + [yPtr] LVARRAY_DEVICE ( std::ptrdiff_t const i ) { PORTABLE_DEVICE_EXPECT_EQ( yPtr[ i ], i ); yPtr[ i ] *= 2; @@ -232,8 +232,8 @@ void testAsyncMemcpyDevice() ArrayView< int, 1, 0, std::ptrdiff_t, BUFFER_TYPE > const yView = y.toView(); 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 ) + RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, y.size() ), + [yView] LVARRAY_DEVICE ( std::ptrdiff_t const i ) { yView[ i ] = -i; } ); @@ -315,8 +315,8 @@ void testAsyncMemcpyDevice() stream.wait_for( e ); 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 ) + RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, y.size() ), + [yPtr] LVARRAY_DEVICE ( std::ptrdiff_t const i ) { PORTABLE_DEVICE_EXPECT_EQ( yPtr[ i ], i ); yPtr[ i ] *= 2; @@ -336,8 +336,8 @@ void testAsyncMemcpyDevice() ArrayView< int, 1, 0, std::ptrdiff_t, BUFFER_TYPE > const yView = y.toView(); 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 ) + RAJA::TypedRangeSegment< std::ptrdiff_t >( 0, y.size() ), + [yView] LVARRAY_DEVICE ( std::ptrdiff_t const i ) { yView[ i ] = -i; } ); diff --git a/unitTests/testStackArray.cpp b/unitTests/testStackArray.cpp index 7c6b6ad7..224f8a09 100644 --- a/unitTests/testStackArray.cpp +++ b/unitTests/testStackArray.cpp @@ -175,7 +175,7 @@ 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() { - if constexpr( std::is_same< POLICY, serialPolicy >::value ) + if constexpr ( std::is_same< POLICY, serialPolicy >::value ) { GTEST_SKIP() << "Nested device lambdas are only supported by a device execution policy."; return; @@ -187,30 +187,30 @@ class StackArrayCaptureTest : public StackArrayTest< typename PERMUTATION_POLICY { dims[ i ] = 8; } 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 ); + { + 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_DEVICE_EXPECT_EQ( array.size( i ), 8 ); } + for( int i = 0; i < NDIM; ++i ) + { PORTABLE_DEVICE_EXPECT_EQ( array.size( i ), 8 ); } - PORTABLE_DEVICE_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_DEVICE_EXPECT_EQ( array.size( 0 ), 2 ); - for( int i = 1; i < NDIM; ++i ) - { PORTABLE_DEVICE_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_DEVICE_EXPECT_EQ( array.size(), array.capacity() / 4 ); + PORTABLE_DEVICE_EXPECT_EQ( array.size(), array.capacity() / 4 ); - forValuesInSliceWithIndices( array.toSlice(), CheckValue() ); - } ); + forValuesInSliceWithIndices( array.toSlice(), CheckValue() ); + } ); } } diff --git a/unitTests/testTensorOpsOneSize.cpp b/unitTests/testTensorOpsOneSize.cpp index 1fe2a5a4..efdbadf5 100644 --- a/unitTests/testTensorOpsOneSize.cpp +++ b/unitTests/testTensorOpsOneSize.cpp @@ -265,7 +265,7 @@ class OneSizeTest : public ::testing::Test tensorOps::scaledCopy< N >( a, b, scale ); \ for( std::ptrdiff_t i = 0; i < N; ++i ) \ { \ - if constexpr( std::is_integral< T >::value ) \ + 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 ); } \ diff --git a/unitTests/testUtils.hpp b/unitTests/testUtils.hpp index b110aa47..164b83ea 100644 --- a/unitTests/testUtils.hpp +++ b/unitTests/testUtils.hpp @@ -125,7 +125,7 @@ LVARRAY_HOST_DEVICE inline bool portableEqual( LHS const & lhs, RHS const & rhs { using LhsValue = std::decay_t< LHS >; using RhsValue = std::decay_t< RHS >; - if constexpr( std::is_floating_point< LhsValue >::value || std::is_floating_point< RhsValue >::value ) + if constexpr ( std::is_floating_point< LhsValue >::value || std::is_floating_point< RhsValue >::value ) { using FloatingPoint = std::common_type_t< LhsValue, RhsValue >; return math::abs( static_cast< FloatingPoint >( lhs ) - static_cast< FloatingPoint >( rhs ) ) <= NumericLimits< FloatingPoint >::epsilon; From e0c5ffdbdfc290fb1bac665c8d0b00f519887c05 Mon Sep 17 00:00:00 2001 From: "Victor A. P. Magri" Date: Fri, 21 Aug 2026 11:36:25 -0400 Subject: [PATCH 6/7] Fix CUDA test comparisons --- unitTests/testUtils.hpp | 20 +------------------- 1 file changed, 1 insertion(+), 19 deletions(-) diff --git a/unitTests/testUtils.hpp b/unitTests/testUtils.hpp index 164b83ea..d26643ba 100644 --- a/unitTests/testUtils.hpp +++ b/unitTests/testUtils.hpp @@ -120,26 +120,8 @@ LAYOUT const & getRAJAViewLayout( RAJA::View< T, LAYOUT > const & view ) #endif } -template< typename LHS, typename RHS > -LVARRAY_HOST_DEVICE inline bool portableEqual( LHS const & lhs, RHS const & rhs ) -{ - using LhsValue = std::decay_t< LHS >; - using RhsValue = std::decay_t< RHS >; - if constexpr ( std::is_floating_point< LhsValue >::value || std::is_floating_point< RhsValue >::value ) - { - using FloatingPoint = std::common_type_t< LhsValue, RhsValue >; - return math::abs( static_cast< FloatingPoint >( lhs ) - static_cast< FloatingPoint >( rhs ) ) <= NumericLimits< FloatingPoint >::epsilon; - } - else - { - return lhs == rhs; - } -} - - #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) -#define PORTABLE_EXPECT_EQ( L, R ) LVARRAY_ERROR_IF( !portableEqual( ( L ), ( R ) ), \ - "Expected " STRINGIZE( L ) " == " STRINGIZE( R ) ) +#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, \ STRINGIZE( L ) " = " << ( L ) << "\n" << STRINGIZE( R ) " = " << ( R ) ); #else From 61cd9b4ab41410cbc81692214f01a475a38038f3 Mon Sep 17 00:00:00 2001 From: "Victor A. P. Magri" Date: Fri, 21 Aug 2026 13:24:11 -0400 Subject: [PATCH 7/7] Use modern Umpire CMake target --- CMakeLists.txt | 8 +++----- cmake/SetupTPL.cmake | 2 +- 2 files changed, 4 insertions(+), 6 deletions(-) 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()