fix: CUDA/HIP unit tests

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 >;

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 );

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
@@ -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

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 )
 

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 )
   {

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 );

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_NEAR_2D( N, M, dstMatrix, result, 100 * NumericLimits< T >::epsilon ); \
+            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 )
           {

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_NEAR_2D( N, M, dstMatrix, result, 100 * NumericLimits< T >::epsilon ); \
+            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 )
           {

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 > >