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/9] 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/9] 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/9] 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/9] 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/9] 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/9] 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/9] 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() From 1c1b7295ec5681829fcbd319b470a46064f7c3fe Mon Sep 17 00:00:00 2001 From: "Victor A. P. Magri" Date: Tue, 25 Aug 2026 09:16:56 -0400 Subject: [PATCH 8/9] Avoid undefined behavior in array views --- src/ArrayOfArraysView.hpp | 27 +++++++++++++++++++++------ src/ArraySlice.hpp | 9 ++++++++- src/ArrayView.hpp | 24 ++++++++++++++++++++---- src/ChaiBuffer.hpp | 2 +- src/indexing.hpp | 22 +++++++++++++++++++--- src/umpireInterface.cpp | 15 +++++++++++++++ 6 files changed, 84 insertions(+), 15 deletions(-) diff --git a/src/ArrayOfArraysView.hpp b/src/ArrayOfArraysView.hpp index 726ec078..59171a56 100644 --- a/src/ArrayOfArraysView.hpp +++ b/src/ArrayOfArraysView.hpp @@ -929,7 +929,10 @@ class ArrayOfArraysView { INDEX_TYPE const curArraySize = sizeOfArray( array ); INDEX_TYPE const curArrayOffset = m_offsets[ array ]; - arrayManipulation::uninitializedShiftUp( &buffer[ curArrayOffset ], curArraySize, capacityIncrease ); + if( curArraySize > 0 ) + { + arrayManipulation::uninitializedShiftUp( &buffer[ curArrayOffset ], curArraySize, capacityIncrease ); + } } } else @@ -944,8 +947,11 @@ class ArrayOfArraysView INDEX_TYPE const curArraySize = sizeOfArray( array ); INDEX_TYPE const curArrayOffset = m_offsets[ array ]; INDEX_TYPE shift = array > i ? capacityIncrease : 0; - arrayManipulation::uninitializedMove( &newBuffer[ curArrayOffset + shift ], curArraySize, &buffer[ curArrayOffset ] ); - arrayManipulation::destroy( &buffer[ curArrayOffset ], curArraySize ); + if( curArraySize > 0 ) + { + arrayManipulation::uninitializedMove( &newBuffer[ curArrayOffset + shift ], curArraySize, &buffer[ curArrayOffset ] ); + arrayManipulation::destroy( &buffer[ curArrayOffset ], curArraySize ); + } } buffer.free(); @@ -968,14 +974,20 @@ class ArrayOfArraysView [this, i, capacityDecrease, arrayOffset, newArraySize, prevArraySize] ( auto & buffer ) { // Delete the values at the end of the array. - arrayManipulation::destroy( &buffer[ arrayOffset + newArraySize ], prevArraySize - newArraySize ); + if( prevArraySize > newArraySize ) + { + arrayManipulation::destroy( &buffer[ arrayOffset + newArraySize ], prevArraySize - newArraySize ); + } // Shift down the values of subsequent arrays. for( INDEX_TYPE array = i + 1; array < m_numArrays; ++array ) { INDEX_TYPE const curArraySize = sizeOfArray( array ); INDEX_TYPE const curArrayOffset = m_offsets[array]; - arrayManipulation::uninitializedShiftDown( &buffer[ curArrayOffset ], curArraySize, capacityDecrease ); + if( curArraySize > 0 ) + { + arrayManipulation::uninitializedShiftDown( &buffer[ curArrayOffset ], curArraySize, capacityDecrease ); + } } }, m_values, buffers ... @@ -1056,7 +1068,10 @@ class ArrayOfArraysView { INDEX_TYPE const offset = m_offsets[ i ]; INDEX_TYPE const arraySize = sizeOfArray( i ); - arrayManipulation::destroy( &buffer[ offset ], arraySize ); + if( arraySize > 0 ) + { + arrayManipulation::destroy( &buffer[ offset ], arraySize ); + } } } }, m_values, buffers ... ); diff --git a/src/ArraySlice.hpp b/src/ArraySlice.hpp index 374979e3..6ec5de62 100644 --- a/src/ArraySlice.hpp +++ b/src/ArraySlice.hpp @@ -292,7 +292,14 @@ class ArraySlice operator[]( INDEX_TYPE const index ) const noexcept { ARRAY_SLICE_CHECK_BOUNDS( index ); - return m_data[ indexing::ConditionalMultiply< USD == 0 >::multiply( index, m_strides[ 0 ] ) ]; + if constexpr( USD == 0 ) + { + return m_data[ index ]; + } + else + { + return m_data[ index * m_strides[ 0 ] ]; + } } /** diff --git a/src/ArrayView.hpp b/src/ArrayView.hpp index d9745d3a..4b6d78f4 100644 --- a/src/ArrayView.hpp +++ b/src/ArrayView.hpp @@ -502,9 +502,18 @@ class ArrayView operator[]( INDEX_TYPE const index ) const & noexcept { ARRAY_SLICE_CHECK_BOUNDS( index ); - return ArraySlice< T, NDIM-1, USD-1, INDEX_TYPE >( data() + indexing::ConditionalMultiply< USD == 0 >::multiply( index, m_strides[ 0 ] ), - m_dims.data + 1, - m_strides.data + 1 ); + if constexpr( USD == 0 ) + { + return ArraySlice< T, NDIM-1, USD-1, INDEX_TYPE >( data() + index, + m_dims.data + 1, + m_strides.data + 1 ); + } + else + { + return ArraySlice< T, NDIM-1, USD-1, INDEX_TYPE >( data() + index * m_strides[ 0 ], + m_dims.data + 1, + m_strides.data + 1 ); + } } /** @@ -531,7 +540,14 @@ class ArrayView operator[]( INDEX_TYPE const index ) const & noexcept { ARRAY_SLICE_CHECK_BOUNDS( index ); - return data()[ indexing::ConditionalMultiply< USD == 0 >::multiply( index, m_strides[ 0 ] ) ]; + if constexpr( USD == 0 ) + { + return data()[ index ]; + } + else + { + return data()[ index * m_strides[ 0 ] ]; + } } /** diff --git a/src/ChaiBuffer.hpp b/src/ChaiBuffer.hpp index f6d4c3d5..5756594c 100644 --- a/src/ChaiBuffer.hpp +++ b/src/ChaiBuffer.hpp @@ -265,7 +265,7 @@ class ChaiBuffer ChaiBuffer( ChaiBuffer< U > const & src ): m_pointer( reinterpret_cast< T * >( src.data() ) ), m_capacity( typeManipulation::convertSize< T, U >( src.capacity() ) ), - m_pointerRecord( &src.pointerRecord() ) + m_pointerRecord( src.data() == nullptr ? nullptr : &src.pointerRecord() ) {} /** diff --git a/src/indexing.hpp b/src/indexing.hpp index d5ea6fe9..ddcf7c2e 100644 --- a/src/indexing.hpp +++ b/src/indexing.hpp @@ -104,7 +104,16 @@ multiplyAll( T const * const LVARRAY_RESTRICT values ) template< int USD, typename INDEX_TYPE, typename INDEX > LVARRAY_HOST_DEVICE inline constexpr INDEX_TYPE getLinearIndex( INDEX_TYPE const * const LVARRAY_RESTRICT strides, INDEX const index ) -{ return ConditionalMultiply< USD == 0 >::multiply( index, strides[ 0 ] ); } +{ + if constexpr( USD == 0 ) + { + return index; + } + else + { + return index * strides[ 0 ]; + } +} /** * @tparam USD The unit stride dimension of strides. @@ -122,8 +131,15 @@ template< int USD, typename INDEX_TYPE, typename INDEX, typename ... REMAINING_I LVARRAY_HOST_DEVICE inline constexpr INDEX_TYPE getLinearIndex( INDEX_TYPE const * const LVARRAY_RESTRICT strides, INDEX const index, REMAINING_INDICES const ... indices ) { - return ConditionalMultiply< USD == 0 >::multiply( index, strides[ 0 ] ) + - getLinearIndex< USD - 1, INDEX_TYPE, REMAINING_INDICES... >( strides + 1, indices ... ); + if constexpr( USD == 0 ) + { + return index + getLinearIndex< USD - 1, INDEX_TYPE, REMAINING_INDICES... >( strides + 1, indices ... ); + } + else + { + return index * strides[ 0 ] + + getLinearIndex< USD - 1, INDEX_TYPE, REMAINING_INDICES... >( strides + 1, indices ... ); + } } /// @return A string representing an empty set of indices. diff --git a/src/umpireInterface.cpp b/src/umpireInterface.cpp index 9219564d..9a36b936 100644 --- a/src/umpireInterface.cpp +++ b/src/umpireInterface.cpp @@ -25,6 +25,11 @@ namespace umpireInterface void copy( void * const dstPointer, void * const srcPointer, std::size_t const size ) { + if( size == 0 ) + { + return; + } + #if defined( LVARRAY_USE_UMPIRE ) umpire::ResourceManager & rm = umpire::ResourceManager::getInstance(); if( rm.hasAllocator( dstPointer ) && rm.hasAllocator( srcPointer ) ) @@ -40,6 +45,11 @@ void copy( void * const dstPointer, void * const srcPointer, std::size_t const s camp::resources::Event copy( void * const dstPointer, void * const srcPointer, camp::resources::Resource & resource, std::size_t const size ) { + if( size == 0 ) + { + return resource.get_event(); + } + #if defined( LVARRAY_USE_UMPIRE ) umpire::ResourceManager & rm = umpire::ResourceManager::getInstance(); @@ -64,6 +74,11 @@ camp::resources::Event copy( void * const dstPointer, void * const srcPointer, void memset( void * const dstPointer, int const val, std::size_t const size ) { + if( size == 0 ) + { + return; + } + #if defined( LVARRAY_USE_UMPIRE ) umpire::ResourceManager & rm = umpire::ResourceManager::getInstance(); if( rm.hasAllocator( dstPointer ) ) From d4e20ee99e34b4db9ce3a5aa4c66394d1955f1d6 Mon Sep 17 00:00:00 2001 From: "Victor A. P. Magri" Date: Wed, 26 Aug 2026 00:15:28 -0400 Subject: [PATCH 9/9] Format constexpr conditions --- src/ArraySlice.hpp | 2 +- src/ArrayView.hpp | 4 ++-- src/indexing.hpp | 4 ++-- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/src/ArraySlice.hpp b/src/ArraySlice.hpp index 6ec5de62..624756d8 100644 --- a/src/ArraySlice.hpp +++ b/src/ArraySlice.hpp @@ -292,7 +292,7 @@ class ArraySlice operator[]( INDEX_TYPE const index ) const noexcept { ARRAY_SLICE_CHECK_BOUNDS( index ); - if constexpr( USD == 0 ) + if constexpr ( USD == 0 ) { return m_data[ index ]; } diff --git a/src/ArrayView.hpp b/src/ArrayView.hpp index 4b6d78f4..e4d548cb 100644 --- a/src/ArrayView.hpp +++ b/src/ArrayView.hpp @@ -502,7 +502,7 @@ class ArrayView operator[]( INDEX_TYPE const index ) const & noexcept { ARRAY_SLICE_CHECK_BOUNDS( index ); - if constexpr( USD == 0 ) + if constexpr ( USD == 0 ) { return ArraySlice< T, NDIM-1, USD-1, INDEX_TYPE >( data() + index, m_dims.data + 1, @@ -540,7 +540,7 @@ class ArrayView operator[]( INDEX_TYPE const index ) const & noexcept { ARRAY_SLICE_CHECK_BOUNDS( index ); - if constexpr( USD == 0 ) + if constexpr ( USD == 0 ) { return data()[ index ]; } diff --git a/src/indexing.hpp b/src/indexing.hpp index ddcf7c2e..1c9ed949 100644 --- a/src/indexing.hpp +++ b/src/indexing.hpp @@ -105,7 +105,7 @@ template< int USD, typename INDEX_TYPE, typename INDEX > LVARRAY_HOST_DEVICE inline constexpr INDEX_TYPE getLinearIndex( INDEX_TYPE const * const LVARRAY_RESTRICT strides, INDEX const index ) { - if constexpr( USD == 0 ) + if constexpr ( USD == 0 ) { return index; } @@ -131,7 +131,7 @@ template< int USD, typename INDEX_TYPE, typename INDEX, typename ... REMAINING_I LVARRAY_HOST_DEVICE inline constexpr INDEX_TYPE getLinearIndex( INDEX_TYPE const * const LVARRAY_RESTRICT strides, INDEX const index, REMAINING_INDICES const ... indices ) { - if constexpr( USD == 0 ) + if constexpr ( USD == 0 ) { return index + getLinearIndex< USD - 1, INDEX_TYPE, REMAINING_INDICES... >( strides + 1, indices ... ); }