Getting HIP device code running on GPU for spheral_cuda_test; revert …

…boost back to 1.74.0; rocmcc error mitigation for building older boost.

cmake/Compilers.cmake

@@ -14,7 +14,7 @@ option(ENABLE_MISSING_INCLUDE_DIR_WARNINGS "show unused parameter warnings" ON)
 set(CXX_WARNING_FLAGS "")
 if (ENABLE_WARNINGS)
   if("${CMAKE_CXX_COMPILER_ID}" STREQUAL "Clang")
-    list(APPEND CXX_WARNING_FLAGS -fdiagnostics-show-option -Wno-unused-command-line-argument -Wno-c++17-extensions)
+    list(APPEND CXX_WARNING_FLAGS -fdiagnostics-show-option -Wno-unused-command-line-argument -Wno-c++17-extensions -Wno-enum-constexpr-conversion)
   endif()
 else()
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -w")

scripts/spack/packages/spheral/package.py

@@ -39,7 +39,7 @@ class Spheral(CachedCMakePackage, CudaPackage, ROCmPackage):
 
     depends_on('[email protected]:', type='build')
 
-    depends_on('boost +system +filesystem -atomic -container -coroutine -chrono -context -date_time -exception -fiber -graph -iostreams -locale -log -math -mpi -program_options -python -random -regex -test -thread -timer -wave +pic', type='build')
+    depends_on('boost@1.74.0 +system +filesystem -atomic -container -coroutine -chrono -context -date_time -exception -fiber -graph -iostreams -locale -log -math -mpi -program_options -python -random -regex -test -thread -timer -wave +pic', type='build')
 
     depends_on('[email protected] +shared +pic', type='build')
 
@@ -139,6 +139,7 @@ def initconfig_mpi_entries(self):
         entries = []
         if "+mpi" in spec:
           entries = super(Spheral, self).initconfig_mpi_entries()
+          # When on cray / flux systems we need to tell CMAKE the mpi flag explicitly
           if "cray-mpich" in spec:
             for e in entries:
                 if 'MPIEXEC_NUMPROC_FLAG' in e:
@@ -152,7 +153,7 @@ def initconfig_hardware_entries(self):
 
         if '+rocm' in spec:
             entries.append(cmake_cache_option("ENABLE_HIP", True))
-            entries.append(cmake_cache_option("ROCM_PATH", spec["hip"].prefix))
+            entries.append(cmake_cache_string("ROCM_PATH", spec["hip"].prefix))
 
         if '+cuda' in spec:
             entries.append(cmake_cache_option("ENABLE_CUDA", True))

scripts/spheral_ats.py

@@ -162,7 +162,7 @@ def main():
         elif any(x in hostname for x in toss_cray_machine_names):
             os.environ['MACHINE_TYPE'] = 'flux00'
             numNodes = numNodes if numNodes else 2
-            timeLimit = timeLimit if timeLimit else 120
+            timeLimit = timeLimit if timeLimit else 60
             #mac_args = [f"--nn={numNodes} --gpus_per_task=1 -n=64 --timelimit={timeLimit}m"]
             #inAllocVars = ["SLURM_JOB_NUM_NODES", "SLURM_NNODES"]
             launch_cmd = f"flux alloc --exclusive -N {numNodes} -t {timeLimit} "

src/CXXTests/DeviceTestLib/DeviceTest.cc

@@ -5,17 +5,19 @@
 namespace Spheral
 {
 
-#ifdef __CUDACC__
+#if defined(__CUDACC__) or defined(__HIPCC__)
 __device__ void add(int a, int b, int *c)
 {
   *c = a + b;
 }
 
 __global__ void launch(int a, int b, int *c)
 {
-  add(a,b,c);
+  *c = a + b;
+  //add(a,b,c);
 }
 
+#if defined(__CUDACC__)
 __host__ int launchCaller(int a, int b)
 {
   int c;
@@ -31,6 +33,25 @@ __host__ int launchCaller(int a, int b)
   cudaFree(d_c);
   return c;
 }
+#endif
+
+#if defined(__HIPCC__)
+__host__ int launchCaller(int a, int b)
+{
+  int c;
+  int *d_c;
+  hipMalloc((void**) &d_c, sizeof(int));
+
+  launch<<<1,1>>>(a,b,d_c);
+  hipError_t err = hipGetLastError();
+  if (err != hipSuccess) 
+      printf("Error: %s\n", hipGetErrorString(err));
+
+  hipMemcpy(&c, d_c, sizeof(int), hipMemcpyDeviceToHost);
+  hipFree(d_c);
+  return c;
+}
+#endif
 
 #else
 int launchCaller(int a, int b)

src/CXXTests/DeviceTestLib/DeviceTest.hh

@@ -1,6 +1,10 @@
 #ifndef __Spheral_DeviceTest_hh__
 #define __Spheral_DeviceTest_hh__
 
+#if defined(SPHERAL_ENABLE_HIP)
+#include <hip/hip_runtime.h>
+#endif
+
 namespace Spheral
 {