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/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..624756d8 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..e4d548cb 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..1c9ed949 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/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<valueString.size(); ++a )
+  std::string valueStringWithSpaces;
+  valueStringWithSpaces.reserve( valueString.size() + numClose );
+  for( char const c : valueString )
   {
-    if( valueString[a] == '}' )
+    if( c == '}' )
     {
-      valueString.insert( a, " " );
-      ++a;
+      valueStringWithSpaces.push_back( ' ' );
     }
+    valueStringWithSpaces.push_back( c );
   }
-  std::istringstream strstream( valueString );
+  std::istringstream strstream( valueStringWithSpaces );
   // this recursively reads the values from the stringstream
   internal::StringToArrayHelper< T, INDEX_TYPE >::Read( array.toSlice(), array.dims(), strstream );
 }
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 ) )
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