Support HIP 4.1.x and breaking HIP_PLATFORM change

CMakeLists.txt

@@ -181,6 +181,8 @@ if(MSVC OR WIN32 OR CYGWIN OR APPLE)
     set(GINKGO_HAVE_HWLOC 0)
 endif()
 
+# We keep using NVCC/HCC for consistency with previous releases even if AMD
+# updated everything to use NVIDIA/AMD in ROCM 4.1
 set(GINKGO_HIP_PLATFORM_NVCC 0)
 set(GINKGO_HIP_PLATFORM_HCC 0)
 
@@ -196,10 +198,12 @@ if(GINKGO_BUILD_HIP)
             "Set and export the environment variable HIP_PLATFORM.")
     endif()
     message(STATUS "HIP platform set to ${GINKGO_HIP_PLATFORM}")
+    set(HIP_PLATFORM_AMD_REGEX "hcc|amd")
+    set(HIP_PLATFORM_NVIDIA_REGEX "nvcc|nvidia")
 
-    if (GINKGO_HIP_PLATFORM STREQUAL "hcc")
+    if (GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_AMD_REGEX}")
         set(GINKGO_HIP_PLATFORM_HCC 1)
-    elseif (GINKGO_HIP_PLATFORM STREQUAL "nvcc")
+    elseif (GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}")
         set(GINKGO_HIP_PLATFORM_NVCC 1)
     endif()
 endif()

INSTALL.md

@@ -220,26 +220,26 @@ imposed by the `HIP` tool suite. The variables are the following:
 
 #### HIP platform detection of AMD and NVIDIA
 By default, Ginkgo uses the output of `/opt/rocm/hip/bin/hipconfig --platform`
-to select the backend. The accepted values are either `hcc` (AMD) or `nvcc`
-(NVIDIA). When on an AMD or NVIDIA system, this should output the correct
-platform by default. When on a system without GPUs, this should output `hcc` by
-default. To change this value, export the environment variable `HIP_PLATFORM`
-like so:
+to select the backend. The accepted values are either `hcc` (`amd` with ROCM >=
+4.1) or `nvcc` (`nvidia` with ROCM >= 4.1). When on an AMD or NVIDIA system,
+this should output the correct platform by default. When on a system without
+GPUs, this should output `hcc` by default. To change this value, export the
+environment variable `HIP_PLATFORM` like so:
 ```bash
-export HIP_PLATFORM=nvcc
+export HIP_PLATFORM=nvcc # or nvidia for ROCM >= 4.1
 ```
 
-When using `HIP_PLATFORM=hcc`, note that two `HIP` compilers can be set, the old
-`hcc` or since ROCm 3.5, `clang`. Ginkgo is only compatible with the `clang`
-based installations. Although this setting should be automatically done, it is
-also possible to manually set the `HIP` compiler to `clang`:
+When using `HIP_PLATFORM=hcc` (or `amd`), note that two `HIP` compilers can be
+set: the old `hcc`, or since ROCm 3.5, `clang`. Ginkgo is only compatible with
+the `clang` based installations. Although this setting should be automatically
+done, it is also possible to manually set the `HIP` compiler to `clang`:
 ```
 export HIP_COMPILER=clang
 ```
 
 #### Setting platform specific compilation flags
-Platform specific compilation flags can be given through the following
-CMake variables:
+Platform specific compilation flags can be given through the following CMake
+variables:
 + `-DGINKGO_HIP_COMPILER_FLAGS=`: compilation flags given to all platforms.
 + `-DGINKGO_HIP_NVCC_COMPILER_FLAGS=`: compilation flags given to NVIDIA platforms.
 + `-DGINKGO_HIP_CLANG_COMPILER_FLAGS=`: compilation flags given to AMD clang compiler.

benchmark/CMakeLists.txt

@@ -4,11 +4,13 @@ if (NOT CMAKE_BUILD_TYPE STREQUAL "Release")
         "will be affected")
 endif()
 
-if (GINKGO_BUILD_CUDA AND GINKGO_BUILD_HIP AND GINKGO_HIP_PLATFORM MATCHES "hcc")
+if (GINKGO_BUILD_CUDA AND GINKGO_BUILD_HIP AND
+    GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_AMD_REGEX}")
     message(FATAL_ERROR "Building the benchmarks for both HIP AMD and CUDA "
         "at the same time is currently not supported. "
         "Disable the benchmark build using `-DGINKGO_BUILD_BENCHMARKS=OFF` "
-        "or use `export HIP_PLATFORM=nvcc` in your build environment instead.")
+        "or use `export HIP_PLATFORM=nvcc` (ROCM <=4.0) or "
+        "`export HIP_PLATFORM=nvidia` (ROCM >= 4.1) in your build environment instead.")
 endif()
 
 function(ginkgo_benchmark_add_tuning_maybe name)

cmake/build_type_helpers.cmake

@@ -91,10 +91,10 @@ foreach(_LANG IN LISTS ENABLED_LANGUAGES ITEMS "HIP")
                 message(STATUS "Skipping ${_LANG}, not supported by build_type.cmake script")
             endif()
             set(${PROJECT_NAME}_${_LANG}_${_TYPE}_SUPPORTED FALSE)
-		        continue()
+            continue()
         endif()
         if(${PROJECT_NAME}_${_LANG}_${_TYPE}_SUPPORTED)
-            if(_LANG STREQUAL "HIP" AND GINKGO_HIP_PLATFORM STREQUAL "nvcc")
+            if(_LANG STREQUAL "HIP" AND GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}")
                 set(CMAKE_${_LANG}_FLAGS_${_TYPE}
                     ${${PROJECT_NAME}_NVCC_${_TYPE}_COMPILER_FLAGS}
                     CACHE STRING "Flags used by the ${_LANG} compiler during ${_TYPE} builds." FORCE

cmake/create_test.cmake

@@ -112,7 +112,7 @@ function(ginkgo_create_hip_test test_name)
     # NOTE: With how HIP works, passing the flags `HIPCC_OPTIONS` etc. here
     # creates a redefinition of all flags. This creates some issues with `nvcc`,
     # but `clang` seems fine with the redefinitions.
-    if (GINKGO_HIP_PLATFORM STREQUAL "nvcc")
+    if (GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}")
         hip_add_executable(${TEST_TARGET_NAME} ${test_name}.hip.cpp
             # If `FindHIP.cmake`, namely `HIP_PARSE_HIPCC_OPTIONS` macro and
             # call gets fixed, uncomment this.

core/CMakeLists.txt

@@ -75,7 +75,7 @@ endif()
 
 # Since we have a public dependency on HIP, this dependency appears
 # here as well
-if (GINKGO_BUILD_HIP AND GINKGO_HIP_PLATFORM MATCHES "hcc")
+if (GINKGO_BUILD_HIP AND GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_AMD_REGEX}")
     list(APPEND GKO_RPATH_ADDITIONS "${HIP_PATH}/lib")
 endif()
 

hip/CMakeLists.txt

@@ -1,4 +1,5 @@
-if (GINKGO_HIP_PLATFORM MATCHES "nvcc" AND CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 9.2)
+if (GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}"
+    AND CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 9.2)
     message(FATAL_ERROR "Ginkgo HIP backend requires CUDA >= 9.2.")
 endif()
 
@@ -95,19 +96,20 @@ execute_process(
             )
 set(GINKGO_HIP_VERSION ${GINKGO_HIP_VERSION} PARENT_SCOPE)
 
-if (GINKGO_HIP_PLATFORM MATCHES "nvcc") # ensure ENV{CUDA_PATH} is set by the user
+if (GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}")
+    # ensure ENV{CUDA_PATH} is set by the user
     if (NOT DEFINED ENV{CUDA_PATH})
         find_path(GINKGO_HIP_DEFAULT_CUDA_PATH "cuda.h" PATH /usr/local/cuda/include NO_DEFAULT_PATH)
         if (NOT GINKGO_HIP_DEFAULT_CUDA_PATH)
-            message(FATAL_ERROR "HIP nvcc backend was requested but CUDA could not be located. "
-                "Set and export the environment variable CUDA_PATH.")
+            message(FATAL_ERROR "HIP nvidia backend was requested but CUDA could not be "
+                "located. Set and export the environment variable CUDA_PATH.")
          endif()
      endif()
 endif()
 
 ## Setup all CMAKE variables to find HIP and its dependencies
 list(APPEND CMAKE_MODULE_PATH "${HIP_PATH}/cmake")
-if (GINKGO_HIP_PLATFORM STREQUAL "hcc")
+if (GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_AMD_REGEX}")
     list(APPEND CMAKE_PREFIX_PATH "${HIP_PATH}/lib/cmake")
 endif()
 list(APPEND CMAKE_PREFIX_PATH
@@ -192,7 +194,7 @@ set(GINKGO_HIP_SOURCES
     stop/residual_norm_kernels.hip.cpp)
 
 set(GINKGO_HIP_NVCC_ARCH "")
-if (GINKGO_HIP_PLATFORM MATCHES "nvcc")
+if (GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}")
     if (NOT CMAKE_CUDA_HOST_COMPILER AND NOT GINKGO_CUDA_DEFAULT_HOST_COMPILER)
         set(CMAKE_CUDA_HOST_COMPILER "${CMAKE_CXX_COMPILER}" CACHE STRING "" FORCE)
     elseif(GINKGO_CUDA_DEFAULT_HOST_COMPILER)
@@ -205,8 +207,9 @@ if (GINKGO_HIP_PLATFORM MATCHES "nvcc")
     # Remove false positive CUDA warnings when calling one<T>() and zero<T>()
     list(APPEND GINKGO_HIP_NVCC_ADDITIONAL_FLAGS "--expt-relaxed-constexpr")
 
-    if (GINKGO_HIP_PLATFROM MATCHES "nvcc" AND CMAKE_CUDA_COMPILER_VERSION
-            MATCHES "9.2" AND CMAKE_CUDA_HOST_COMPILER MATCHES ".*clang.*" )
+    if (GINKGO_HIP_PLATFROM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}"
+            AND CMAKE_CUDA_COMPILER_VERSION MATCHES "9.2"
+            AND CMAKE_CUDA_HOST_COMPILER MATCHES ".*clang.*" )
         ginkgo_extract_clang_version(${CMAKE_CUDA_HOST_COMPILER} GINKGO_CUDA_HOST_CLANG_VERSION)
 
         if (GINKGO_CUDA_HOST_CLANG_VERSION MATCHES "5\.0.*")
@@ -255,7 +258,7 @@ target_link_libraries(ginkgo_hip PRIVATE roc::hipblas roc::hipsparse hip::hipran
 
 target_compile_options(ginkgo_hip PRIVATE $<$<COMPILE_LANGUAGE:CXX>:${GINKGO_COMPILER_FLAGS}>)
 
-if(GINKGO_HIP_PLATFORM MATCHES "hcc")
+if(GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_AMD_REGEX}")
     find_package(hip REQUIRED)
     # To save a bit of pain, we directly link against the `library` instead of
     # linking against the target.
@@ -266,14 +269,14 @@ if(GINKGO_HIP_PLATFORM MATCHES "hcc")
         get_target_property(HIP_LIBAMDHIP64_LIBRARIES hip::amdhip64 IMPORTED_LOCATION_RELEASE)
     endif()
     target_link_libraries(ginkgo_hip PUBLIC ${HIP_LIBAMDHIP64_LIBRARIES})
-elseif(GINKGO_HIP_PLATFORM MATCHES "nvcc")
+elseif(GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}")
     find_package(CUDA 9.2 REQUIRED)
     target_link_libraries(ginkgo_hip PUBLIC ${CUDA_LIBRARIES})
 endif()
 
 # Try to find everything in /opt/rocm/lib first.
 set(GKO_HIP_RPATH "${ROCM_PATH}/lib" )
-if (GINKGO_HIP_PLATFORM MATCHES "nvcc")
+if (GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}")
     list(GET CUDA_LIBRARIES 0 CUDA_FIRST_LIB)
     get_filename_component(GKO_CUDA_LIBDIR "${CUDA_FIRST_LIB}" DIRECTORY)
     list(APPEND GKO_HIP_RPATH "${GKO_CUDA_LIBDIR}")

hip/test/base/CMakeLists.txt

@@ -5,7 +5,7 @@ endif()
 ginkgo_create_hip_test(lin_op)
 ginkgo_create_hip_test(math)
 # Only hcc needs the libraries. nvcc only requires the headers.
-if (GINKGO_HIP_PLATFORM MATCHES "hcc")
+if (GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_AMD_REGEX}")
     ginkgo_create_hip_test(exception_helpers roc::hipblas roc::hipsparse hip::hiprand roc::rocrand)
 else()
     ginkgo_create_hip_test(exception_helpers)

third_party/CMakeLists.txt

@@ -1,4 +1,5 @@
-if(GINKGO_BUILD_CUDA OR (GINKGO_BUILD_HIP AND GINKGO_HIP_PLATFORM STREQUAL "nvcc"))
+if(GINKGO_BUILD_CUDA OR (GINKGO_BUILD_HIP AND
+        GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}"))
     enable_language(CUDA)
     if (GINKGO_USE_EXTERNAL_CAS)
         include(CudaArchitectureSelector RESULT_VARIABLE GINKGO_CAS_FILE)