fix some cuda issues...more to come

Author: rrsettgast

cmake/blt

@@ -0,0 +1,83 @@
+#################################################################################
+# Generated host-config - Edit at own risk!
+#################################################################################
+#--------------------------------------------------------------------------------
+# SYS_TYPE: toss_4_x86_64_ib
+# Compiler Spec: gcc@=13.3.1
+# CMake executable path: /usr/tce/backend/installations/linux-rhel8-x86_64/gcc-10.3.1/cmake-3.26.3-nz532rvfpaf5lf74zxmplgiobuhol7lu/bin/cmake
+#--------------------------------------------------------------------------------
+
+#--------------------------------------------------------------------------------
+# Compilers
+#--------------------------------------------------------------------------------
+
+set(CMAKE_C_COMPILER "/usr/tce/packages/gcc/gcc-13.3.1-magic/bin/gcc" CACHE PATH "")
+
+set(CMAKE_CXX_COMPILER "/usr/tce/packages/gcc/gcc-13.3.1-magic/bin/g++" CACHE PATH "")
+
+set(CMAKE_CXX_FLAGS_RELEASE "-O3 -DNDEBUG" CACHE STRING "")
+
+set(CMAKE_CXX_FLAGS_RELWITHDEBINFO "-O2 -g -DNDEBUG" CACHE STRING "")
+
+set(CMAKE_CXX_FLAGS_DEBUG "-g" CACHE STRING "")
+
+#--------------------------------------------------------------------------------
+# CMake Standard
+#--------------------------------------------------------------------------------
+
+set(BLT_CXX_STD "c++17" CACHE STRING "")
+
+
+#--------------------------------------------------------------------------------
+# Cuda
+#--------------------------------------------------------------------------------
+
+set(ENABLE_CUDA ON CACHE BOOL "")
+
+set(CMAKE_CUDA_STANDARD "17" CACHE PATH "")
+
+set(CUDA_TOOLKIT_ROOT_DIR "/usr/tce/packages/cuda/cuda-12.9.1" CACHE PATH "")
+
+set(CMAKE_CUDA_COMPILER "${CUDA_TOOLKIT_ROOT_DIR}/bin/nvcc" CACHE PATH "")
+
+set(CMAKE_CUDA_ARCHITECTURES "90" CACHE STRING "")
+
+set(CMAKE_CUDA_FLAGS "-restrict --expt-extended-lambda -Werror cross-execution-space-call,reorder,deprecated-declarations -arch sm_90" CACHE STRING "")
+
+set(CMAKE_CUDA_FLAGS_RELWITHDEBINFO "-g -lineinfo ${CMAKE_CUDA_FLAGS_RELEASE}" CACHE STRING "")
+
+set(CMAKE_CUDA_FLAGS_DEBUG "-g -G -O0 -Xcompiler -O0" CACHE STRING "")
+
+
+#--------------------------------------------------------------------------------
+# System Math Libraries
+#--------------------------------------------------------------------------------
+
+set(ENABLE_MKL ON CACHE BOOL "")
+
+set(MKL_INCLUDE_DIRS "/usr/tce/packages/mkl/mkl-2022.1.0/include" CACHE PATH "")
+
+set(MKL_LIBRARIES /usr/tce/packages/mkl/mkl-2022.1.0/mkl/2022.1.0/lib/intel64/libmkl_intel_lp64.so
+                  /usr/tce/packages/mkl/mkl-2022.1.0/mkl/2022.1.0/lib/intel64/libmkl_gnu_thread.so
+                  /usr/tce/packages/mkl/mkl-2022.1.0/mkl/2022.1.0/lib/intel64/libmkl_core.so
+                  /usr/tce/backend/installations/linux-rhel8-x86_64/gcc-13.3.1/llvm-19.1.3-gy2lu5xbi4csr2k47emlajzfs5mlsd4g/bin/../lib/x86_64-unknown-linux-gnu/libomp.so
+                  /lib64/libpthread.so
+                  /lib64/libm.so
+                  /lib64/libdl.so CACHE STRING "")
+
+
+#--------------------------------------------------------------------------------
+# Documentation
+#--------------------------------------------------------------------------------
+
+set(ENABLE_DOCS OFF CACHE BOOL "")
+
+set(ENABLE_DOXYGEN OFF CACHE BOOL "")
+
+set(ENABLE_SPHINX OFF CACHE BOOL "")
+
+#--------------------------------------------------------------------------------
+# Development tools
+#--------------------------------------------------------------------------------
+
+set(ENABLE_UNCRUSTIFY OFF CACHE BOOL "")

hostconfigs/LLNL/[email protected]@12.9.1.cmake

@@ -26,15 +26,16 @@ set( hpcReact_sources)
 find_package(LAPACK REQUIRED)
 find_package(BLAS REQUIRED)
 
-
-
-#target_link_libraries(MyExecutable PRIVATE LAPACK::LAPACK)
 
-set( hpcReack_dependencies
+set( hpcReact_dependencies
      LAPACK::LAPACK
      BLAS::BLAS
      )
 
+if( ENABLE_CUDA )
+  list( APPEND hpcReact_dependencies cuda )
+endif()
+
 blt_add_library( NAME             hpcReact
 #                 SOURCES          ${hpcReact_sources}
                  HEADERS          ${hpcReact_headers}

src/CMakeLists.txt

@@ -39,6 +39,7 @@ template< typename T, int DIM0 >
 struct CArrayWrapper< T, DIM0 >
 {
   // default constructor
+  HPCREACT_HOST_DEVICE
   constexpr CArrayWrapper() = default;
 
   /**
@@ -53,6 +54,7 @@ struct CArrayWrapper< T, DIM0 >
    *
    * @note No runtime bounds checking is performed on the initializer size.
    */
+  HPCREACT_HOST_DEVICE
   constexpr CArrayWrapper( std::initializer_list< T > init )
   {
     // static_assert(init.size() == DIM0, "Size mismatch"); // needs c++20
@@ -66,6 +68,7 @@ struct CArrayWrapper< T, DIM0 >
    * @brief Copy constructor.
    * @param src The source CArrayWrapper to copy from.
    */
+  HPCREACT_HOST_DEVICE
   constexpr CArrayWrapper( CArrayWrapper const & src )
   {
     for( std::size_t i = 0; i < DIM0; i++ )
@@ -79,27 +82,31 @@ struct CArrayWrapper< T, DIM0 >
    * @param dim The index (must be in range [0, DIM0)).
    * @return Reference to the element at the specified index.
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T & operator()( int const dim ) { return data[dim]; }
 
   /**
    * @brief Read-only access to an element by index (const overload).
    * @param dim The index (must be in range [0, DIM0)).
    * @return Const reference to the element at the specified index.
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T const & operator()( int const dim ) const { return data[dim]; }
 
   /**
    * @brief Subscript operator for read/write access.
    * @param dim The index (must be in range [0, DIM0)).
    * @return Reference to the element at the specified index.
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T & operator[]( int const dim ) { return data[dim]; }
 
   /**
    * @brief Subscript operator for read-only access (const overload).
    * @param dim The index (must be in range [0, DIM0)).
    * @return Const reference to the element at the specified index.
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T const & operator[]( int const dim ) const { return data[dim]; }
 
   /// The underlying 1D C-style array.
@@ -120,12 +127,14 @@ template< typename T, int DIM0, int DIM1 >
 struct CArrayWrapper< T, DIM0, DIM1 >
 {
   // default constructor
+  HPCREACT_HOST_DEVICE
   constexpr CArrayWrapper() = default;
 
   /**
    * @brief Copy constructor.
    * @param src The source CArrayWrapper to copy from.
    */
+  HPCREACT_HOST_DEVICE
   constexpr CArrayWrapper( CArrayWrapper const & src )
   {
     for( std::size_t i = 0; i < DIM0; i++ )
@@ -150,6 +159,7 @@ struct CArrayWrapper< T, DIM0, DIM1 >
    *
    * @note No runtime bounds checking is performed on the initializer dimensions.
    */
+  HPCREACT_HOST_DEVICE
   constexpr CArrayWrapper( std::initializer_list< std::initializer_list< T > > init )
   {
     // static_assert(init.size() == DIM0, "Size mismatch"); // needs c++20
@@ -172,6 +182,7 @@ struct CArrayWrapper< T, DIM0, DIM1 >
    * @param dim1 Index in the second dimension (range [0, DIM1)).
    * @return Reference to the element at the specified 2D location.
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T & operator()( int const dim0, int const dim1 )
   {
     return data[dim0][dim1];
@@ -183,6 +194,7 @@ struct CArrayWrapper< T, DIM0, DIM1 >
    * @param dim1 Index in the second dimension (range [0, DIM1)).
    * @return Const reference to the element at the specified 2D location.
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T const & operator()( int const dim0, int const dim1 ) const
   {
     return data[dim0][dim1];
@@ -195,6 +207,7 @@ struct CArrayWrapper< T, DIM0, DIM1 >
    *
    * This allows usage like `obj[dim0][dim1]`.
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T       ( & operator[]( int const dim0 ))[DIM1]
   {
     return data[dim0];
@@ -205,6 +218,7 @@ struct CArrayWrapper< T, DIM0, DIM1 >
    * @param dim0 The row index (range [0, DIM0)).
    * @return Const reference to an array of type T[DIM1].
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T const (&operator[]( int const dim0 ) const)[DIM1]
   {
     return data[dim0];
@@ -229,6 +243,7 @@ template< typename T, int DIM0, int DIM1, int DIM2 >
 struct CArrayWrapper< T, DIM0, DIM1, DIM2 >
 {
   // default constructor
+  HPCREACT_HOST_DEVICE
   constexpr CArrayWrapper() = default;
 
   /**
@@ -254,6 +269,7 @@ struct CArrayWrapper< T, DIM0, DIM1, DIM2 >
    * @note This constructor does not perform size validation. Incorrect initializer sizes
    *       may lead to undefined behavior.
    */
+  HPCREACT_HOST_DEVICE
   constexpr CArrayWrapper( std::initializer_list< std::initializer_list< std::initializer_list< T > > > init )
   {
     // static_assert(init.size() == DIM0, "Size mismatch"); // needs c++20
@@ -286,6 +302,7 @@ struct CArrayWrapper< T, DIM0, DIM1, DIM2 >
    * @note Currently, this function incorrectly indexes data[dim0][dim1], missing dim2.
    *       It should be `data[dim0][dim1][dim2]`. Please correct if intended.
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T & operator()( int const dim0, int const dim1, int const dim2 )
   {
     // NOTE: This looks like a bug in your original code. Should be data[dim0][dim1][dim2].
@@ -299,6 +316,7 @@ struct CArrayWrapper< T, DIM0, DIM1, DIM2 >
    * @param dim2 Index in the third dimension (range [0, DIM2)).
    * @return Const reference to the element at the specified 3D location.
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T const & operator()( int const dim0, int const dim1, int const dim2 ) const
   {
     // NOTE: Same potential bug as above. Should be data[dim0][dim1][dim2].
@@ -312,6 +330,7 @@ struct CArrayWrapper< T, DIM0, DIM1, DIM2 >
    *
    * This allows usage like `obj[dim0][dim1][dim2]`.
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T ( & operator[]( int const dim0 ))[DIM1][DIM2]
   {
     return data[dim0];
@@ -322,6 +341,7 @@ struct CArrayWrapper< T, DIM0, DIM1, DIM2 >
    * @param dim0 The slice index (range [0, DIM0)).
    * @return Const reference to an array of type T[DIM1][DIM2].
    */
+  HPCREACT_HOST_DEVICE
   constexpr HPCREACT_HOST_DEVICE inline T const (&operator[]( int const dim0 ) const)[DIM1][DIM2]
   {
     return data[dim0];

src/common/CArrayWrapper.hpp

@@ -19,6 +19,7 @@ namespace hpcReact
 {
 
 template< typename REAL_TYPE, int N >
+HPCREACT_HOST_DEVICE
 bool isPositiveDefinite( REAL_TYPE const (&A)[N][N] )
 {
   REAL_TYPE temp[N][N];
@@ -52,6 +53,7 @@ bool isPositiveDefinite( REAL_TYPE const (&A)[N][N] )
 
 
 template< typename REAL_TYPE, int N >
+HPCREACT_HOST_DEVICE
 void solveNxN_Cholesky( REAL_TYPE const (&A)[N][N], REAL_TYPE const (&b)[N], REAL_TYPE (& x)[N] )
 {
   REAL_TYPE L[N][N] = {{0}};
@@ -96,6 +98,7 @@ void solveNxN_Cholesky( REAL_TYPE const (&A)[N][N], REAL_TYPE const (&b)[N], REA
 
 
 template< typename REAL_TYPE, int N >
+HPCREACT_HOST_DEVICE
 void solveNxN_Cholesky( symmetricMatrix< REAL_TYPE, int, N > const & A,
                         REAL_TYPE const (&b)[N],
                         REAL_TYPE (& x)[N] )

src/common/DirectSystemSolve.hpp

@@ -13,18 +13,23 @@
 
 
 
-#if defined( __CUDACC__ ) || defined( __HIPCC__ )
+#if defined( __CUDACC__ )
 #define HPCREACT_USE_DEVICE
-#define HPCREACT_HOST_DEVICE __host__ __device__
-#else
-#define HPCREACT_HOST_DEVICE
+#define HPCREACT_USE_CUDA
+#elif defined( __HIPCC__ )
+#define HPCREACT_USE_DEVICE
+#define HPCREACT_USE_HIP
 #endif
 
 
 #if defined( HPCREACT_USE_DEVICE )
 #define HPCREACT_GLOBAL __global__
+#define HPCREACT_DEVICE __device__
+#define HPCREACT_HOST_DEVICE __host__ __device__
 #else
 #define HPCREACT_GLOBAL
+#define HPCREACT_DEVICE
+#define HPCREACT_HOST_DEVICE
 #endif
 
 /// This macro is used to ignore warnings that that a variable is

src/common/macros.hpp

@@ -13,24 +13,26 @@
 
 #include "common/macros.hpp"
 
+#include <cstdio>
 #include <utility>
+#include <iostream>
 
 
 namespace hpcReact
 {
 #if defined(HPCREACT_USE_DEVICE)
   #if defined(HPCREACT_USE_CUDA)
-#define deviceMalloc( PTR, BYTES ) cudaMalloc( PTR, BYTES );
-#define deviceMallocManaged( PTR, BYTES ) cudaMallocManaged( PTR, BYTES );
-#define deviceDeviceSynchronize() cudaDeviceSynchronize();
-#define deviceMemCpy( DST, SRC, BYTES, KIND ) cudaMemcpy( DST, SRC, BYTES, KIND );
-#define deviceFree( PTR ) cudaFree( PTR );
+    #define deviceMalloc( PTR, BYTES ) cudaMalloc( PTR, BYTES );
+    #define deviceMallocManaged( PTR, BYTES ) cudaMallocManaged( PTR, BYTES );
+    #define deviceDeviceSynchronize() cudaDeviceSynchronize();
+    #define deviceMemCpy( DST, SRC, BYTES, KIND ) cudaMemcpy( DST, SRC, BYTES, KIND );
+    #define deviceFree( PTR ) cudaFree( PTR );
   #elif defined(HPCREACT_USE_HIP)
-#define deviceMalloc( PTR, BYTES ) hipMalloc( PTR, BYTES );
-#define deviceMallocManaged( PTR, BYTES ) hipMallocManaged( PTR, BYTES );
-#define deviceDeviceSynchronize() hipDeviceSynchronize();
-#define deviceMemCpy( DST, SRC, BYTES, KIND ) hipMemcpy( DST, SRC, BYTES, KIND );
-#define deviceFree( PTR ) hipFree( PTR );
+    #define deviceMalloc( PTR, BYTES ) hipMalloc( PTR, BYTES );
+    #define deviceMallocManaged( PTR, BYTES ) hipMallocManaged( PTR, BYTES );
+    #define deviceDeviceSynchronize() hipDeviceSynchronize();
+    #define deviceMemCpy( DST, SRC, BYTES, KIND ) hipMemcpy( DST, SRC, BYTES, KIND );
+    #define deviceFree( PTR ) hipFree( PTR );
   #endif
 #endif
 
@@ -121,13 +123,20 @@ HPCREACT_GLOBAL void genericKernel( LAMBDA func, DATA_TYPE * const data )
 template< typename DATA_TYPE, typename LAMBDA >
 void genericKernelWrapper( int const N, DATA_TYPE * const hostData, LAMBDA && func )
 {
-
 #if defined(HPCREACT_USE_DEVICE)
   DATA_TYPE * deviceData;
   deviceMalloc( &deviceData, N * sizeof(DATA_TYPE) );
   deviceMemCpy( deviceData, hostData, N * sizeof(DATA_TYPE), cudaMemcpyHostToDevice );
   genericKernel <<< 1, 1 >>> ( std::forward< LAMBDA >( func ), deviceData );
+
+  cudaError_t e = cudaGetLastError();
+  if (e != cudaSuccess) { fprintf(stderr, "launch error: %s\n", cudaGetErrorString(e)); abort(); }
+
   deviceDeviceSynchronize();
+
+  e = cudaGetLastError();
+  if (e != cudaSuccess) { fprintf(stderr, "post-sync error: %s\n", cudaGetErrorString(e)); abort(); }
+
   deviceMemCpy( hostData, deviceData, N * sizeof(DATA_TYPE), cudaMemcpyDeviceToHost );
   deviceFree( deviceData );
 #else

src/common/pmpl.hpp

@@ -2,8 +2,13 @@
 set( testSourceFiles
      testDirectSystemSolve.cpp )
 
+
 set( dependencyList hpcReact gtest )
 
+if( ENABLE_CUDA )
+    list( APPEND dependencyList cuda )
+endif()
+
 # Add gtest C++ based tests
 foreach(test ${testSourceFiles})
     get_filename_component( test_name ${test} NAME_WE )