From 6594c6ad51dc1ad8b37d02fbb3392c3ec066f0ff Mon Sep 17 00:00:00 2001
From: "Victor A. P. Magri" <paludettomag1@llnl.gov>
Date: Thu, 27 Aug 2026 23:47:23 -0400
Subject: [PATCH 1/6] Update wait_for calls

---
 unitTests/testMemcpy.cpp | 16 ++++++++--------
 1 file changed, 8 insertions(+), 8 deletions(-)

diff --git a/unitTests/testMemcpy.cpp b/unitTests/testMemcpy.cpp
index 0e44243d..3f045461 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 )
   {
@@ -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 319341c5c3b6b391af97e1c7fe491da324d57fa2 Mon Sep 17 00:00:00 2001
From: "Victor A. P. Magri" <paludettomag1@llnl.gov>
Date: Fri, 28 Aug 2026 03:16:19 -0400
Subject: [PATCH 2/6] Fix device asinh precision

---
 src/math.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/math.hpp b/src/math.hpp
index 0d358b32..0d5e6a09 100644
--- a/src/math.hpp
+++ b/src/math.hpp
@@ -1075,7 +1075,7 @@ LVARRAY_HOST_DEVICE LVARRAY_FORCE_INLINE
 float asinh( float const x )
 {
 #if defined(LVARRAY_DEVICE_COMPILE)
-  return ::asinhf( x );
+  return static_cast< float >( ::asinh( static_cast< double >( x ) ) );
 #else
   return std::asinh( x );
 #endif

From c4d950c923be8ab2a63426dcef091dcf09cd1079 Mon Sep 17 00:00:00 2001
From: "Victor A. P. Magri" <paludettomag1@llnl.gov>
Date: Fri, 28 Aug 2026 03:43:44 -0400
Subject: [PATCH 3/6] Use tolerance for scaled add tests

---
 unitTests/testTensorOpsCommon.hpp    | 17 +++++++++++++++++
 unitTests/testTensorOpsTwoSizes.hpp  |  3 ++-
 unitTests/testTensorOpsTwoSizes1.cpp |  3 ++-
 3 files changed, 21 insertions(+), 2 deletions(-)

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/testTensorOpsTwoSizes.hpp b/unitTests/testTensorOpsTwoSizes.hpp
index 5492b2b5..cda38449 100644
--- a/unitTests/testTensorOpsTwoSizes.hpp
+++ b/unitTests/testTensorOpsTwoSizes.hpp
@@ -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, result[ N - 1 ][ M - 1 ] * epsilon ); \
 
           #define _TEST_PERMS( dstMatrix, srcMatrix0, srcMatrix1, srcMatrix2, srcMatrix3 ) \
             _TEST( dstMatrix, srcMatrix0 ); \
@@ -701,6 +701,7 @@ class TwoSizesTest : public ::testing::Test
           fill( matrixA_local, matrixSeed );
 
           T const scale = T( 3.14 );
+          T const epsilon = NumericLimitsNC< T >{}.epsilon;
           T result[ N ][ M ];
           for( std::ptrdiff_t i = 0; i < N; ++i )
           {
diff --git a/unitTests/testTensorOpsTwoSizes1.cpp b/unitTests/testTensorOpsTwoSizes1.cpp
index 96ac793c..1384391b 100644
--- a/unitTests/testTensorOpsTwoSizes1.cpp
+++ b/unitTests/testTensorOpsTwoSizes1.cpp
@@ -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, result[ N - 1 ][ M - 1 ] * epsilon ); \
 
           #define _TEST_PERMS( dstMatrix, srcMatrix0, srcMatrix1, srcMatrix2, srcMatrix3 ) \
             _TEST( dstMatrix, srcMatrix0 ); \
@@ -702,6 +702,7 @@ class TwoSizesTest : public ::testing::Test
           fill( matrixA_local, matrixSeed );
 
           T const scale = T( 3.14 );
+          T const epsilon = NumericLimitsNC< T >{}.epsilon;
           T result[ N ][ M ];
           for( std::ptrdiff_t i = 0; i < N; ++i )
           {

From 4276b9f60651fec1bb6aaa9289b2329d8d528c11 Mon Sep 17 00:00:00 2001
From: "Victor A. P. Magri" <paludettomag1@llnl.gov>
Date: Fri, 28 Aug 2026 12:37:02 -0400
Subject: [PATCH 4/6] Fix HIP portability after develop merge

---
 src/Macros.hpp             |  4 ++--
 src/bufferManipulation.hpp |  2 +-
 src/math.hpp               | 34 +++++++++++++++---------------
 unitTests/testMemcpy.cpp   |  2 ++
 unitTests/testUtils.hpp    | 42 ++++++++++++++++++++++++++++++++------
 5 files changed, 58 insertions(+), 26 deletions(-)

diff --git a/src/Macros.hpp b/src/Macros.hpp
index e5ec9ff0..77f29c8c 100644
--- a/src/Macros.hpp
+++ b/src/Macros.hpp
@@ -708,7 +708,7 @@ struct RAJAHelper< parallelHostPolicy >
 
 #endif
 
-#if defined(LVARRAY_USE_CUDA)
+#if defined(LVARRAY_USE_CUDA) && defined(RAJA_CUDA_ACTIVE)
 
 template< unsigned long THREADS_PER_BLOCK >
 using parallelDevicePolicy = RAJA::cuda_exec< THREADS_PER_BLOCK >;
@@ -720,7 +720,7 @@ struct RAJAHelper< RAJA::cuda_exec< N > >
   using AtomicPolicy = RAJA::cuda_atomic;
 };
 
-#elif defined(LVARRAY_USE_HIP)
+#elif defined(LVARRAY_USE_HIP) && defined(RAJA_HIP_ACTIVE)
 
 template< unsigned long THREADS_PER_BLOCK >
 using parallelDevicePolicy = RAJA::hip_exec< THREADS_PER_BLOCK >;
diff --git a/src/bufferManipulation.hpp b/src/bufferManipulation.hpp
index 83e5a00e..a02c2eac 100644
--- a/src/bufferManipulation.hpp
+++ b/src/bufferManipulation.hpp
@@ -191,7 +191,7 @@ void free( BUFFER & buf, std::ptrdiff_t const size )
 
   check( buf, size );
 
-  if( !std::is_trivially_destructible< T >::value )
+  if constexpr( !std::is_trivially_destructible< T >::value )
   {
     buf.move( MemorySpace::host, true );
     arrayManipulation::destroy( buf.data(), size );
diff --git a/src/math.hpp b/src/math.hpp
index 0d5e6a09..6c195954 100644
--- a/src/math.hpp
+++ b/src/math.hpp
@@ -20,7 +20,7 @@
 #include <cmath>
 #include <type_traits>
 
-#if defined( LVARRAY_USE_CUDA )
+#if defined( LVARRAY_USE_CUDA ) && defined( LVARRAY_DECORATE )
   #include <cuda_fp16.h>
 #endif
 
@@ -101,7 +101,7 @@ LVARRAY_HOST_DEVICE inline constexpr
 T lessThan( T const x, T const y )
 { return __hlt( x, y ); }
 
-#if defined( LVARRAY_USE_CUDA )
+#if defined( LVARRAY_USE_CUDA ) && defined( LVARRAY_DECORATE )
 /**
  * @brief Convert @p u to @c __half.
  * @tparam U The type to convert from.
@@ -207,7 +207,7 @@ __half getSecond( __half2 const x )
 
 #endif
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 /**
  * @return 1 if @p x is less than @p y, else 0.
  * @param x The first value.
@@ -319,7 +319,7 @@ max( T const a, T const b )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc max( T, T )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -368,7 +368,7 @@ min( T const a, T const b )
 #endif
 }
 
-#if defined( LVARRAY_USE_CUDA )
+#if defined( LVARRAY_USE_CUDA ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc min( T, T )
 LVARRAY_DEVICE
@@ -414,7 +414,7 @@ T abs( T const x )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc abs( T )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -485,7 +485,7 @@ double sqrt( T const x )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc sqrt( float )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -527,7 +527,7 @@ double invSqrt( T const x )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc invSqrt( float )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -576,7 +576,7 @@ double sin( T const theta )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc sin( float )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -618,7 +618,7 @@ double cos( T const theta )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc cos( float )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -682,7 +682,7 @@ void sincos( T const theta, double & sinTheta, double & cosTheta )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc sincos( float, float &, float & )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -730,7 +730,7 @@ double tan( T const theta )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc tan( float )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -875,7 +875,7 @@ double asin( T const x )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc asin( float )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -917,7 +917,7 @@ double acos( T const x )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc acos( float )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -960,7 +960,7 @@ double atan2( T const y, T const x )
 #endif
 }
 
-#if defined( LVARRAY_USE_CUDA )
+#if defined( LVARRAY_USE_CUDA ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc atan2( float, float )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -1009,7 +1009,7 @@ double exp( T const x )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc exp( float )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
@@ -1051,7 +1051,7 @@ double log( T const x )
 #endif
 }
 
-#if defined( LVARRAY_USE_DEVICE )
+#if defined( LVARRAY_USE_DEVICE ) && defined( LVARRAY_DECORATE )
 
 /// @copydoc log( float )
 LVARRAY_DEVICE LVARRAY_FORCE_INLINE
diff --git a/unitTests/testMemcpy.cpp b/unitTests/testMemcpy.cpp
index c82d569a..0d216680 100644
--- a/unitTests/testMemcpy.cpp
+++ b/unitTests/testMemcpy.cpp
@@ -324,6 +324,7 @@ void testAsyncMemcpyDevice()
 
   e = memcpy< 0, 0 >( stream, x, {}, y.toViewConst(), {} );
   stream.wait_for( e );
+  stream.wait();
 
   for( std::ptrdiff_t i = 0; i < x.size(); ++i )
   {
@@ -344,6 +345,7 @@ void testAsyncMemcpyDevice()
 
   e = memcpy< 0, 0 >( stream, x, {}, y.toViewConst(), {} );
   stream.wait_for( e );
+  stream.wait();
 
   for( std::ptrdiff_t i = 0; i < x.size(); ++i )
   {
diff --git a/unitTests/testUtils.hpp b/unitTests/testUtils.hpp
index d26643ba..baff57cb 100644
--- a/unitTests/testUtils.hpp
+++ b/unitTests/testUtils.hpp
@@ -62,7 +62,7 @@ struct RAJAHelper< parallelHostPolicy >
 
 #endif
 
-#if defined(LVARRAY_USE_CUDA)
+#if defined(LVARRAY_USE_CUDA) && defined(RAJA_CUDA_ACTIVE)
 
 template< unsigned long THREADS_PER_BLOCK >
 using parallelDevicePolicy = RAJA::cuda_exec< THREADS_PER_BLOCK >;
@@ -76,7 +76,7 @@ struct RAJAHelper< RAJA::policy::cuda::cuda_exec_explicit< X, Y, C, BLOCK_SIZE,
   static constexpr MemorySpace space = MemorySpace::cuda;
 };
 
-#elif defined(LVARRAY_USE_HIP)
+#elif defined(LVARRAY_USE_HIP) && defined(RAJA_HIP_ACTIVE)
 
 template< unsigned long THREADS_PER_BLOCK >
 using parallelDevicePolicy = RAJA::hip_exec< THREADS_PER_BLOCK >;
@@ -120,7 +120,27 @@ LAYOUT const & getRAJAViewLayout( RAJA::View< T, LAYOUT > const & view )
 #endif
 }
 
-#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)
+#if defined(__HIPCC__)
+// Clang HIP parses device lambdas during the host pass as well. GoogleTest's
+// host-only assertion macros therefore cannot be used in those lambdas. Keep
+// the check valid in both passes; a failing device-side check traps the kernel.
+#define PORTABLE_EXPECT_EQ( L, R ) \
+  do \
+  { \
+    if( !(( L ) == ( R )) ) \
+    { \
+      __builtin_trap(); \
+    } \
+  } while( false )
+#define PORTABLE_EXPECT_NEAR( L, R, EPSILON ) \
+  do \
+  { \
+    if( math::abs( ( L ) -( R ) ) > ( EPSILON ) ) \
+    { \
+      __builtin_trap(); \
+    } \
+  } while( false );
+#elif 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, \
                                                                        STRINGIZE( L ) " = " << ( L ) << "\n" << STRINGIZE( R ) " = " << ( R ) );
@@ -130,9 +150,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.
+// A device-only lambda is compiled by HIP in a host pass as well. Keep this
+// assertion independent of GoogleTest and avoid device-side printf output.
+#if defined(__HIPCC__)
+#define PORTABLE_DEVICE_EXPECT_EQ( L, R ) \
+  do \
+  { \
+    if( !( ( L ) == ( R ) ) ) \
+    { \
+      __builtin_trap(); \
+    } \
+  } while( false )
+#else
 #define PORTABLE_DEVICE_EXPECT_EQ( L, R ) \
   do \
   { \
@@ -141,6 +170,7 @@ LAYOUT const & getRAJAViewLayout( RAJA::View< T, LAYOUT > const & view )
       assert( false && "Device assertion failed" ); \
     } \
   } while( false )
+#endif
 
 // Comparator that compares a std::pair by it's first object.
 template< class A, class B, class COMP=std::less< A > >

From b147fc1baa41e38ff3f4100bd8e150904ada7b31 Mon Sep 17 00:00:00 2001
From: "Victor A. P. Magri" <paludettomag1@llnl.gov>
Date: Fri, 28 Aug 2026 23:51:15 -0400
Subject: [PATCH 5/6] Fix HIP unit tests

---
 unitTests/CMakeLists.txt           | 8 ++++++++
 unitTests/testTensorOpsInverse.hpp | 7 ++++++-
 2 files changed, 14 insertions(+), 1 deletion(-)

diff --git a/unitTests/CMakeLists.txt b/unitTests/CMakeLists.txt
index 3f5bb961..85a67b3e 100644
--- a/unitTests/CMakeLists.txt
+++ b/unitTests/CMakeLists.txt
@@ -140,6 +140,14 @@ blt_add_executable( NAME testTensorOps
 
 target_include_directories( testTensorOps PUBLIC ${CMAKE_CURRENT_LIST_DIR}/../src )
 
+# ROCm 7.2/amdclang can generate a non-retiring kernel for this exhaustive
+# single-kernel test on gfx10/gfx11 when all tensor operations are inlined.
+# Keep the workaround local to the test target so normal LvArray kernels retain
+# their usual optimization settings.
+if( ENABLE_HIP AND CMAKE_HIP_ARCHITECTURES MATCHES "(^|;)gfx(10|11)" )
+    target_compile_options( testTensorOps PRIVATE -fno-inline )
+endif()
+
 blt_add_test( NAME testTensorOps
               COMMAND testTensorOps )
 
diff --git a/unitTests/testTensorOpsInverse.hpp b/unitTests/testTensorOpsInverse.hpp
index 9edfa950..efcf5c78 100644
--- a/unitTests/testTensorOpsInverse.hpp
+++ b/unitTests/testTensorOpsInverse.hpp
@@ -305,7 +305,12 @@ class InverseTest : public ::testing::Test
     // The bounds for this specific check need to be increased a lot for XL. About 100x for
     // 2x2 even more for 3x3. I'm not sure why, especially since the check below passes.
     #if !defined( __ibmxl__ ) || defined( __CUDA_ARCH__ )
-    PORTABLE_EXPECT_NEAR( det, tensorOps::determinant< M >( source ), scale * epsilon );
+    // The absolute roundoff in a determinant scales with the matrix entries raised to
+    // the matrix dimension. The previous scale * epsilon bound was too strict for
+    // independent 3x3 evaluation orders on HIP.
+    double determinantScale = 1.0;
+    for( int i = 0; i < M; ++i ) determinantScale *= scale;
+    PORTABLE_EXPECT_NEAR( det, tensorOps::determinant< M >( source ), determinantScale * epsilon );
     #endif
 
     PORTABLE_EXPECT_NEAR( 1.0 / det, tensorOps::determinant< M >( inverse ), scale * epsilon );

From 0f6e9eff520f0d664d3169c3f8b08514f613ffeb Mon Sep 17 00:00:00 2001
From: "Victor A. P. Magri" <paludettomag1@llnl.gov>
Date: Sat, 29 Aug 2026 02:59:06 -0400
Subject: [PATCH 6/6] Fix formatting

---
 src/bufferManipulation.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/bufferManipulation.hpp b/src/bufferManipulation.hpp
index a02c2eac..e998fcc9 100644
--- a/src/bufferManipulation.hpp
+++ b/src/bufferManipulation.hpp
@@ -191,7 +191,7 @@ void free( BUFFER & buf, std::ptrdiff_t const size )
 
   check( buf, size );
 
-  if constexpr( !std::is_trivially_destructible< T >::value )
+  if constexpr ( !std::is_trivially_destructible< T >::value )
   {
     buf.move( MemorySpace::host, true );
     arrayManipulation::destroy( buf.data(), size );