feature: add RAJA kernel launches and basic CUDA support

.gitignore

@@ -13,5 +13,6 @@
 *.orig
 __pycache__/
 view
+*.cache*
 /_serac_build_and_test*
 build-linux-*-*-*

CMakeLists.txt

@@ -25,7 +25,7 @@ if (NOT SERAC_STYLE_CI_ONLY)
     set(ENABLE_MPI ON CACHE BOOL "")
 
     if (NOT MPI_C_COMPILER OR NOT MPI_CXX_COMPILER)
-        message(FATAL_ERROR 
+        message(FATAL_ERROR
                 "Serac requires MPI.  It is required to provide the MPI C/C++ "
                 "compiler wrappers via the CMake variables, "
                 "MPI_C_COMPILER and MPI_CXX_COMPILER.")
@@ -146,7 +146,7 @@ endif()
 if (SERAC_STYLE_CI_ONLY)
     # Exit processing the rest of the build in style only build to avoid any possible
     # CMake configuration issues outside of just enabling `make style`. This build,
-    # is not capable of building Serac and should not be advertised. It is for CI 
+    # is not capable of building Serac and should not be advertised. It is for CI
     # purposes only.
     return()
 endif()
@@ -212,7 +212,7 @@ if (SERAC_ENABLE_CODEVELOP)
             )
 endif()
 
-install(EXPORT serac-targets 
+install(EXPORT serac-targets
         NAMESPACE serac::
         DESTINATION lib/cmake
         )

cmake/SeracMacros.cmake

@@ -8,7 +8,7 @@
 #------------------------------------------------------------------------------
 # Adds code checks for all cpp/hpp files recursively under the current directory
 # that regex match INCLUDES and excludes any files that regex match EXCLUDES
-# 
+#
 # This creates the following parent build targets:
 #  check - Runs a non file changing style check and CppCheck
 #  style - In-place code formatting
@@ -71,7 +71,7 @@ macro(serac_add_code_checks)
     blt_add_clang_tidy_target(NAME              ${arg_PREFIX}_guidelines_check_tests
                               CHECKS            "clang-analyzer-*,clang-analyzer-cplusplus*,cppcoreguidelines-*,-cppcoreguidelines-avoid-magic-numbers"
                               SRC_FILES         ${_test_sources})
-                                  
+
     if (ENABLE_COVERAGE)
         blt_add_code_coverage_target(NAME   ${arg_PREFIX}_coverage
                                      RUNNER ${CMAKE_MAKE_PROGRAM} test

cmake/thirdparty/SetupSeracThirdParty.cmake

@@ -23,7 +23,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
         # Manually set includes as system includes
         foreach(_target cuda_runtime cuda)
             get_target_property(_dirs ${_target} INTERFACE_INCLUDE_DIRECTORIES)
-            set_property(TARGET ${_target} 
+            set_property(TARGET ${_target}
                          APPEND PROPERTY INTERFACE_SYSTEM_INCLUDE_DIRECTORIES
                          "${_dirs}")
         endforeach()
@@ -114,7 +114,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
 
     # Manually set includes as system includes
     get_target_property(_dirs conduit::conduit INTERFACE_INCLUDE_DIRECTORIES)
-    set_property(TARGET conduit::conduit 
+    set_property(TARGET conduit::conduit
                  APPEND PROPERTY INTERFACE_SYSTEM_INCLUDE_DIRECTORIES
                  "${_dirs}")
 
@@ -125,7 +125,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
         # Note: Sundials is currently only used via MFEM and MFEM's target contains it's information
         serac_assert_is_directory(DIR_VARIABLE SUNDIALS_DIR)
         set(SERAC_USE_SUNDIALS ON CACHE BOOL "")
-        
+
         # Note: MFEM sets SUNDIALS_FOUND itself
         if (NOT SERAC_ENABLE_CODEVELOP)
             set(SUNDIALS_FOUND TRUE)
@@ -143,7 +143,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
         serac_assert_is_directory(DIR_VARIABLE PETSC_DIR)
         # NOTE: PETSc is built and used through MFEM
         set(SERAC_USE_PETSC ON CACHE BOOL "")
-        
+
         # Note: MFEM *does not* set PETSC_FOUND itself, likely because we skip petsc build tests
         set(PETSC_FOUND TRUE)
     else()
@@ -159,7 +159,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
         serac_assert_is_directory(DIR_VARIABLE SLEPC_DIR)
         # NOTE: SLEPc is built and used through MFEM
         set(SERAC_USE_SLEPC ON CACHE BOOL "")
-        
+
         # Note: MFEM sets SLEPC_FOUND itself
         if (NOT SERAC_ENABLE_CODEVELOP)
             set(SLEPC_FOUND TRUE)
@@ -196,7 +196,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
 
         #### Store Data that MFEM clears
         set(tpls_to_save ADIAK AMGX AXOM CALIPER CAMP CONDUIT HDF5
-                         HYPRE LUA METIS MFEM NETCDF PARMETIS PETSC RAJA 
+                         HYPRE LUA METIS MFEM NETCDF PARMETIS PETSC RAJA
                          SLEPC SUPERLU_DIST STRUMPACK SUNDIALS TRIBOL
                          UMPIRE)
         foreach(_tpl ${tpls_to_save})
@@ -311,7 +311,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
         else()
             add_subdirectory(${PROJECT_SOURCE_DIR}/mfem  ${CMAKE_BINARY_DIR}/mfem)
         endif()
- 
+
         set(MFEM_FOUND TRUE CACHE BOOL "" FORCE)
 
         # Patch the mfem target with the correct include directories
@@ -447,7 +447,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
         else()
             set(TRIBOL_FOUND OFF)
         endif()
-        
+
         message(STATUS "Tribol support is " ${TRIBOL_FOUND})
     else()
         set(ENABLE_FORTRAN OFF CACHE BOOL "" FORCE)
@@ -466,7 +466,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
         endif()
 
         add_subdirectory(${tribol_repo_dir}  ${CMAKE_BINARY_DIR}/tribol)
-        
+
         target_include_directories(redecomp PUBLIC
             $<BUILD_INTERFACE:${tribol_repo_dir}/src>
         )
@@ -475,7 +475,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
             $<BUILD_INTERFACE:${CMAKE_BINARY_DIR}/tribol/include>
             $<INSTALL_INTERFACE:include>
         )
-        
+
         set(TRIBOL_FOUND TRUE CACHE BOOL "" FORCE)
         set(ENABLE_FORTRAN ON CACHE BOOL "" FORCE)
     endif()
@@ -525,7 +525,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
     if(STRUMPACK_DIR)
         list(GET MPI_C_LIBRARIES 0 _first_mpi_lib)
         get_filename_component(_mpi_lib_dir ${_first_mpi_lib} DIRECTORY)
-    
+
         foreach(_target ${_mfem_targets})
             if(TARGET ${_target})
                 message(STATUS "Adding MPI link directory to target [${_target}]")
@@ -591,7 +591,7 @@ if (NOT SERAC_THIRD_PARTY_LIBRARIES_FOUND)
                 message(FATAL_ERROR "Serac+Caliper+CUDA requires CMake > 3.17.")
             else()
                 find_package(CUDAToolkit REQUIRED)
-            endif() 
+            endif()
         endif()
 
         find_dependency(caliper REQUIRED PATHS "${CALIPER_DIR}")

scripts/llnl/common_build_functions.py

@@ -6,7 +6,7 @@
 """
  file: common_build_functions.py
 
- description: 
+ description:
   helpers for installing src and tpls on llnl lc systems.
 
 """
@@ -374,7 +374,7 @@ def build_and_test_host_configs(prefix, timestamp, use_generated_host_configs, r
 
     test_root =  get_build_and_test_root(prefix, timestamp)
     os.mkdir(test_root)
-    write_build_info(pjoin(test_root,"info.json")) 
+    write_build_info(pjoin(test_root,"info.json"))
     ok  = []
     bad = []
     for host_config in host_configs:
@@ -417,7 +417,7 @@ def build_and_test_host_configs(prefix, timestamp, use_generated_host_configs, r
 def set_group_and_perms(directory):
     """
     Sets the proper group and access permissions of given input
-    directory. 
+    directory.
     """
 
     skip = True

src/serac/infrastructure/accelerator.hpp

@@ -12,7 +12,7 @@
  */
 
 #pragma once
-
+#include "RAJA/RAJA.hpp"
 #if defined(__CUDACC__)
 #define SERAC_HOST_DEVICE __host__ __device__
 #define SERAC_HOST __host__
@@ -72,6 +72,26 @@ enum class ExecutionSpace
   Dynamic  // Corresponds to execution that can "legally" happen on either the host or device
 };
 
+#ifdef SERAC_USE_CUDA_KERNEL_EVALUATION
+
+/// @brief Alias for parallel threads policy on GPU
+using threads_x     = RAJA::LoopPolicy<RAJA::cuda_thread_x_direct>;
+using teams_e       = RAJA::LoopPolicy<RAJA::cuda_block_x_direct>;
+using launch_policy = RAJA::LaunchPolicy<RAJA::cuda_launch_t<false>>;
+using forall_policy = RAJA::cuda_exec<128>;
+
+#else
+
+/// @brief Alias for parallel threads policy on GPU.
+using threads_x = RAJA::LoopPolicy<RAJA::seq_exec>;
+/// @brief Alias for number of teams for GPU kernel launches.
+using teams_e = RAJA::LoopPolicy<RAJA::seq_exec>;
+/// @brief Alias for GPU kernel launch policy.
+using launch_policy = RAJA::LaunchPolicy<RAJA::seq_launch_t>;
+using forall_policy = RAJA::seq_exec;
+
+#endif
+
 /**
  * @brief The default execution space for serac builds
  */
@@ -88,6 +108,13 @@ struct execution_to_memory {
   static constexpr axom::MemorySpace value = axom::MemorySpace::Dynamic;
 };
 
+/// @brief This helper is needed to suppress -Werror compilation errors caused by the
+/// explicit captures in the main execution lambdas.
+template <typename... T>
+SERAC_HOST_DEVICE void suppress_unused_capture_warnings(T...)
+{
+}
+
 #ifdef SERAC_USE_UMPIRE
 /// @overload
 template <>

src/serac/infrastructure/debug_print.hpp

@@ -131,16 +131,18 @@ void printCUDAMemUsage()
 {
   int deviceCount = 0;
   cudaGetDeviceCount(&deviceCount);
-  int i = 0;
-  cudaSetDevice(i);
+  for (int i = 0; i < deviceCount; ++i) {
+    cudaSetDevice(i);
 
-  size_t freeBytes, totalBytes;
-  cudaMemGetInfo(&freeBytes, &totalBytes);
-  size_t usedBytes = totalBytes - freeBytes;
+    size_t freeBytes, totalBytes;
+    cudaMemGetInfo(&freeBytes, &totalBytes);
+    size_t usedBytes = totalBytes - freeBytes;
 
-  std::cout << "Device Number: " << i << std::endl;
-  std::cout << " Total Memory (MB): " << (totalBytes / 1024.0 / 1024.0) << std::endl;
-  std::cout << " Free Memory (MB): " << (freeBytes / 1024.0 / 1024.0) << std::endl;
-  std::cout << " Used Memory (MB): " << (usedBytes / 1024.0 / 1024.0) << std::endl;
+    std::cout << "Device Number: " << i << std::endl;
+    std::cout << " Total Memory (MB): " << (totalBytes / 1024.0 / 1024.0) << std::endl;
+    std::cout << " Free Memory (MB): " << (freeBytes / 1024.0 / 1024.0) << std::endl;
+    std::cout << " Used Memory (MB): " << (usedBytes / 1024.0 / 1024.0) << std::endl;
+  }
+  cudaSetDevice(0);
 }
 #endif