Test PR to see how merging https://github.com/ADACS-Australia/bifrost/tree/hipify would go

config/cuda.m4

@@ -8,13 +8,7 @@ AC_DEFUN([AX_CHECK_CUDA],
               [with_cuda_home=/usr/local/cuda])
   AC_SUBST(CUDA_HOME, $with_cuda_home)
 
-  AC_ARG_ENABLE([cuda],
-                [AS_HELP_STRING([--disable-cuda],
-                                [disable cuda support (default=no)])],
-                [enable_cuda=no],
-                [enable_cuda=yes])
-
-  AC_SUBST([HAVE_CUDA], [0])
+  AC_SUBST([HAVE_CUDA], [1])
   AC_SUBST([CUDA_VERSION], [0])
   AC_SUBST([CUDA_HAVE_CXX20], [0])
   AC_SUBST([CUDA_HAVE_CXX17], [0])
@@ -23,7 +17,7 @@ AC_DEFUN([AX_CHECK_CUDA],
   AC_SUBST([GPU_MIN_ARCH], [0])
   AC_SUBST([GPU_MAX_ARCH], [0])
   AC_SUBST([GPU_SHAREDMEM], [0])
-  AC_SUBST([GPU_PASCAL_MANAGEDMEM], [0])
+  AC_SUBST([GPU_MANAGEDMEM], [0])
   AC_SUBST([GPU_EXP_PINNED_ALLOC], [1])
   if test "$enable_cuda" != "no"; then
     AC_SUBST([HAVE_CUDA], [1])
@@ -40,7 +34,7 @@ AC_DEFUN([AX_CHECK_CUDA],
     LDFLAGS_save="$LDFLAGS"
     LIBS_save="$LIBS"
 
-    ac_compile='$NVCC -c $NVCCFLAGS conftest.$ac_ext >&5'
+    ac_compile='$NVCC -c $HIPCCFLAGS conftest.$ac_ext >&5'
     AC_COMPILE_IFELSE([
       AC_LANG_PROGRAM([[
           #include <cuda.h>
@@ -51,9 +45,9 @@ AC_DEFUN([AX_CHECK_CUDA],
 
     if test "$HAVE_CUDA" = "1"; then
       LDFLAGS="-L$CUDA_HOME/lib64 -L$CUDA_HOME/lib"
-      LIBS="$LIBS -lcuda -lcudart"
+      LIBS="-lcuda -lcudart"
 
-      ac_link='$NVCC -o conftest$ac_exeext $NVCCFLAGS $LDFLAGS $LIBS conftest.$ac_ext >&5'
+      ac_link='$NVCC -o conftest$ac_exeext $HIPCCFLAGS $LDFLAGS $LIBS conftest.$ac_ext >&5'
       AC_LINK_IFELSE([
         AC_LANG_PROGRAM([[
             #include <cuda.h>
@@ -105,7 +99,7 @@ AC_DEFUN([AX_CHECK_CUDA],
                               [flags to pass to NVCC (default='-O3 -Xcompiler "-Wall"')])],
               [],
               [with_nvcc_flags='-O3 -Xcompiler "-Wall"'])
-  AC_SUBST(NVCCFLAGS, $with_nvcc_flags)
+  HIPCCFLAGS="$with_nvcc_flags $HIPCCFLAGS"
 
   AC_ARG_WITH([stream_model],
               [AS_HELP_STRING([--with-stream-model],
@@ -119,11 +113,11 @@ AC_DEFUN([AX_CHECK_CUDA],
     dsm_supported=$( ${NVCC} -h | ${GREP} -Po -e "--default-stream" )
     if test "$dsm_supported" = "--default-stream"; then
       if test "$with_stream_model" = "per-thread"; then
-        NVCCFLAGS="$NVCCFLAGS -default-stream per-thread"
+        HIPCCFLAGS="$HIPCCFLAGS -default-stream per-thread"
         AC_MSG_RESULT([yes, using 'per-thread'])
       else
         if test "$with_stream_model" = "legacy"; then
-          NVCCFLAGS="$NVCCFLAGS -default-stream legacy"
+          HIPCCFLAGS="$HIPCCFLAGS -default-stream legacy"
           AC_MSG_RESULT([yes, using 'legacy'])
         else
           AC_MSG_ERROR(Invalid CUDA stream model: '$with_stream_model')
@@ -137,9 +131,9 @@ AC_DEFUN([AX_CHECK_CUDA],
   if test "$HAVE_CUDA" = "1"; then
     CPPFLAGS="$CPPFLAGS -DBF_CUDA_ENABLED=1"
     CXXFLAGS="$CXXFLAGS -DBF_CUDA_ENABLED=1"
-    NVCCFLAGS="$NVCCFLAGS -DBF_CUDA_ENABLED=1"
+    HIPCCFLAGS="$HIPCCFLAGS -DBF_CUDA_ENABLED=1"
     LDFLAGS="$LDFLAGS -L$CUDA_HOME/lib64 -L$CUDA_HOME/lib"
-    LIBS="$LIBS -lcuda -lcudart -lnvrtc -lcublas -lcudadevrt -L. -lcufft_static_pruned -lculibos -lnvToolsExt"
+    LIBS="$LIBS -lcuda -lcudart -lnvrtc -lcublas -lcudadevrt -L. -lnvToolsExt -Wl,--no-as-needed -lcufft_static_pruned -lculibos"
   fi
 
   AC_ARG_WITH([gpu_archs],
@@ -291,10 +285,10 @@ AC_DEFUN([AX_CHECK_CUDA],
     AC_MSG_CHECKING([for Pascal-style CUDA managed memory])
     cm_invalid=$( echo $GPU_ARCHS | ${SED} -e 's/\b[[1-5]][[0-9]]\b/PRE/g;' )
     if ! echo $cm_invalid | ${GREP} -q PRE; then
-      AC_SUBST([GPU_PASCAL_MANAGEDMEM], [1])
+      AC_SUBST([GPU_MANAGEDMEM], [1])
       AC_MSG_RESULT([yes])
     else
-      AC_SUBST([GPU_PASCAL_MANAGEDMEM], [0])
+      AC_SUBST([GPU_MANAGEDMEM], [0])
       AC_MSG_RESULT([no])
     fi
 

config/hip.m4

@@ -0,0 +1,121 @@
+AC_DEFUN([AX_CHECK_HIP],
+[
+    AC_PROVIDE([AX_CHECK_HIP])
+
+    AC_ARG_ENABLE(
+        [gpu],
+        [AS_HELP_STRING(
+            [--enable-gpu],
+            [enable GPU support (default=no)])
+        ],
+        [AC_SUBST([ENABLE_GPU], [1])],
+        [AC_SUBST([ENABLE_GPU], [0])]
+    )
+
+    AC_MSG_CHECKING([if gpu is enabled])
+    AC_MSG_RESULT([$ENABLE_GPU])
+
+    AC_SUBST([GPU_SHAREDMEM], 0)
+    AC_ARG_WITH([shared_mem],
+        [AS_HELP_STRING([--with-shared-mem=N],
+                        [default GPU shared memory per block in bytes (default=detect)])],
+        [AC_SUBST([GPU_SHAREDMEM], [$withval])],
+        [with_shared_mem='auto'])
+    AC_MSG_NOTICE([-with-shared-mem=$GPU_SHAREDMEM])
+
+    AS_IF([test "x$ENABLE_GPU" = "x1"], [
+        AC_PATH_PROG(HIPCONFIG, hipconfig, no)
+        AS_IF([test "x$HIPCONFIG" = "xno"], [
+            AC_MSG_ERROR("could not find hipconfig in path")
+        ])
+
+        AC_PATH_PROG(HIPCC, hipcc, no)
+        AS_IF([test "x$HIPCC" = "xno"], [
+            AC_MSG_ERROR("could not find hipcc in path")
+        ])
+
+        AC_MSG_CHECKING([for GPU platform])
+        AC_SUBST([GPU_PLATFORM], [`hipconfig -P`])
+        AC_MSG_RESULT([$GPU_PLATFORM])
+
+        AC_MSG_CHECKING([for hip path])
+        AC_SUBST([HIP_PATH], [`hipconfig -p`])
+        AC_MSG_RESULT([$HIP_PATH])
+
+        AC_MSG_CHECKING([for rocm path])
+        AC_SUBST([ROCM_PATH], [`hipconfig -R`])
+        AC_MSG_RESULT([$ROCM_PATH])
+
+        AC_MSG_CHECKING([for hipcc C++ config])
+        AC_SUBST([HIP_CPPCONF], [`hipconfig -C`])
+        AC_MSG_RESULT([$HIP_CPPCONF])
+    ])
+
+    AC_SUBST([HIPCCFLAGS])
+
+    AS_IF([test "x$GPU_PLATFORM" = "xnvidia"], [
+        CXXFLAGS="$CXXFLAGS -DBF_CUDA_ENABLED=1 $HIP_CPPCONF"
+        HIPCCFLAGS="$HIPCCFLAGS --std=c++17 -DBF_CUDA_ENABLED=1 $HIP_CPPCONF"
+        LDFLAGS="$LDFLAGS -L$HIP_PATH/lib -L$ROCM_PATH/lib"
+        LIBS="$LIBS -lhipfft -lhipblas"
+    ])
+
+    AS_IF([test "x$GPU_PLATFORM" = "xamd"], [
+        CXXFLAGS="$CXXFLAGS -DBF_CUDA_ENABLED=1 $HIP_CPPCONF "
+        HIPCCFLAGS="$HIPCCFLAGS -std=c++17 -Wall -O3 -DBF_CUDA_ENABLED=1 $HIP_CPPCONF"
+        LDFLAGS="$LDFLAGS -L$HIP_PATH/lib -L$ROCM_PATH/lib"
+        LIBS="$LIBS -lamdhip64 -lhipfft -lhipblas -lhiprtc"
+
+        # AMD Constants
+        AC_SUBST([GPU_MANAGEDMEM], [1])
+        AC_SUBST([GPU_MIN_ARCH], [0])
+        AC_SUBST([GPU_MAX_ARCH], [0])
+        AC_SUBST([GPU_EXP_PINNED_ALLOC], [0])
+        AC_SUBST([CUDA_VERSION], [0])
+
+        AS_IF([test "x$with_shared_mem" = "xauto"], [
+            AC_MSG_CHECKING([GPU shared memory using automatic method])
+
+            AC_RUN_IFELSE([
+                AC_LANG_PROGRAM([[
+                    #include <algorithm>
+                    #include <fstream>
+                    #include <iostream>
+                    #include <limits>
+                    #include <hip/hip_runtime.h>
+                ]], [[
+                    int count {};
+                    auto hiperr = hipGetDeviceCount(&count);
+                    if (hiperr != hipSuccess) {
+                        std::cerr << "Error detecting devices" << std::endl;
+                        return 1;
+                    }
+                    if (count == 0) {
+                        std::cerr << "No devices detected" << std::endl;
+                        return 1;
+                    }
+                    size_t mem {std::numeric_limits<size_t>::max()};
+                    for (int device = 0; device < count; ++device) {
+                        hipDeviceProp_t prop;
+                        hiperr = hipGetDeviceProperties(&prop, device);
+                        if (hiperr != hipSuccess) {
+                            std::cerr << "Failed to query shared memory for device " << device << std::endl;
+                            return 1;
+                        }
+                        mem = std::min(mem, prop.sharedMemPerBlock);
+                    }
+                    std::ofstream fd;
+                    fd.open("confmem.out");
+                    fd << mem;
+                    fd.close();
+                    return 0;
+                ]])
+            ], [
+                AC_SUBST([GPU_SHAREDMEM], [$(cat confmem.out)])
+                AC_MSG_RESULT([$GPU_SHAREDMEM bytes])
+            ], [
+                AC_MSG_FAILURE([automatic shared memory detection failed (error code: $?)])
+            ])
+        ])
+    ])
+])