Modernize CMake setup for HIP

.gitlab-ci.yml

@@ -106,7 +106,6 @@ build/cuda101/nompi/clang/cuda_wo_omp/release/shared:
   variables:
     CXX_COMPILER: "clang++"
     BUILD_CUDA: "ON"
-    BUILD_HIP: "ON"
     BUILD_HWLOC: "OFF"
     BUILD_TYPE: "Release"
     CUDA_ARCH: 35
@@ -124,7 +123,6 @@ build/cuda101/openmpi/gcc/all/debug/shared:
     BUILD_CUDA: "ON"
     BUILD_MPI: "ON"
     MPI_AS_ROOT: "ON"
-    BUILD_HIP: "ON"
     BUILD_TYPE: "Debug"
     BUILD_PAPI_SDE: "ON"
     CUDA_ARCH: 35
@@ -139,7 +137,6 @@ build/cuda101/nompi/clang/all/release/static:
     CXX_COMPILER: "clang++"
     BUILD_OMP: "ON"
     BUILD_CUDA: "ON"
-    BUILD_HIP: "ON"
     BUILD_TYPE: "Release"
     BUILD_SHARED_LIBS: "OFF"
     CUDA_ARCH: 35
@@ -190,7 +187,6 @@ build/cuda102/nompi/gcc/all/debug/shared:
   variables:
     BUILD_OMP: "ON"
     BUILD_CUDA: "ON"
-    BUILD_HIP: "ON"
     BUILD_TYPE: "Debug"
     FAST_TESTS: "ON"
     BUILD_HWLOC: "OFF"
@@ -207,7 +203,6 @@ build/cuda102/nompi/clang/all/release/static:
     CXX_COMPILER: "clang++"
     BUILD_OMP: "ON"
     BUILD_CUDA: "ON"
-    BUILD_HIP: "ON"
     BUILD_TYPE: "Release"
     BUILD_SHARED_LIBS: "OFF"
     CUDA_ARCH: 35
@@ -425,38 +420,38 @@ build/amd/nompi/clang/rocm45/debug/shared:
     BUILD_HIP: "ON"
     BUILD_TYPE: "Debug"
 
-# ROCm 5.0.2 and friends
-build/amd/nompi/gcc/rocm502/debug/static:
+# ROCm 5.1.4 and friends
+build/amd/nompi/gcc/rocm514/debug/static:
   extends:
     - .build_and_test_template
     - .default_variables
     - .full_test_condition
-    - .use_gko-rocm502-nompi-gnu11-llvm11
+    - .use_gko-rocm514-nompi-gnu11-llvm11
   variables:
     BUILD_OMP: "ON"
     BUILD_HIP: "ON"
     BUILD_TYPE: "Debug"
     BUILD_SHARED_LIBS: "OFF"
 
-build/amd/nompi/clang/rocm502/release/shared:
+build/amd/nompi/clang/rocm514/release/shared:
   extends:
     - .build_and_test_template
     - .default_variables
     - .quick_test_condition
-    - .use_gko-rocm502-nompi-gnu11-llvm11
+    - .use_gko-rocm514-nompi-gnu11-llvm11
   variables:
     CXX_COMPILER: "clang++"
     BUILD_OMP: "ON"
     BUILD_HIP: "ON"
     BUILD_TYPE: "Release"
 
 # without omp
-build/amd/nompi/gcc/rocm502_wo_omp/release/shared:
+build/amd/nompi/gcc/rocm514_wo_omp/release/shared:
   extends:
     - .build_and_test_template
     - .default_variables
     - .full_test_condition
-    - .use_gko-rocm502-nompi-gnu11-llvm11
+    - .use_gko-rocm514-nompi-gnu11-llvm11
   variables:
     BUILD_OMP: "OFF"
     BUILD_MPI: "OFF"
@@ -686,7 +681,6 @@ warnings:
   variables:
     BUILD_OMP: "ON"
     BUILD_CUDA: "ON"
-    BUILD_HIP: "ON"
     CXX_FLAGS: "-Werror=pedantic;-pedantic-errors"
   allow_failure: yes
 
@@ -701,7 +695,6 @@ no-circular-deps:
   variables:
     BUILD_OMP: "ON"
     BUILD_CUDA: "ON"
-    BUILD_HIP: "ON"
     EXTRA_CMAKE_FLAGS: '-DGINKGO_CHECK_CIRCULAR_DEPS=on'
   allow_failure: no
 
@@ -729,7 +722,6 @@ clang-tidy:
   variables:
     BUILD_OMP: "ON"
     BUILD_CUDA: "ON"
-    BUILD_HIP: "ON"
     EXTRA_CMAKE_FLAGS: '-DGINKGO_WITH_CLANG_TIDY=ON'
   allow_failure: yes
 

.gitlab/image.yml

@@ -74,8 +74,8 @@
     - private_ci
     - amd-gpu
 
-.use_gko-rocm502-nompi-gnu11-llvm11:
-  image: ginkgohub/rocm:502-openmpi-gnu11-llvm11
+.use_gko-rocm514-nompi-gnu11-llvm11:
+  image: ginkgohub/rocm:514-openmpi-gnu11-llvm11
   tags:
     - private_ci
     - amd-gpu

.gitlab/scripts.yml

@@ -106,7 +106,9 @@
     - popd
     - if [ -n "${SYCL_DEVICE_TYPE}" ]; then unset SYCL_DEVICE_TYPE; fi
     - if [ -n "${SYCL_DEVICE_FILTER}" ]; then unset SYCL_DEVICE_FILTER; fi
-    - PKG_CONFIG_PATH=${INSTALL_PREFIX}/lib/pkgconfig:$PKG_CONFIG_PATH LD_LIBRARY_PATH=${INSTALL_PREFIX}/lib:$LD_LIBRARY_PATH ninja test_pkgconfig
+    - if [[ "${BUILD_SHARED_LIBS}" == "ON" || "${BUILD_HIP}" != "ON" ]]; then
+      PKG_CONFIG_PATH=${INSTALL_PREFIX}/lib/pkgconfig:$PKG_CONFIG_PATH LD_LIBRARY_PATH=${INSTALL_PREFIX}/lib:$LD_LIBRARY_PATH ninja test_pkgconfig;
+      fi
   dependencies: []
 
 
@@ -136,7 +138,9 @@
     - pushd test/test_install
     - ninja install
     - popd
-    - PKG_CONFIG_PATH=${INSTALL_PREFIX}/lib/pkgconfig:$PKG_CONFIG_PATH LD_LIBRARY_PATH=${INSTALL_PREFIX}/lib:$LD_LIBRARY_PATH ninja test_pkgconfig
+    - if [[ "${BUILD_SHARED_LIBS}" == "ON" || "${BUILD_HIP}" != "ON" ]]; then
+      PKG_CONFIG_PATH=${INSTALL_PREFIX}/lib/pkgconfig:$PKG_CONFIG_PATH LD_LIBRARY_PATH=${INSTALL_PREFIX}/lib:$LD_LIBRARY_PATH ninja test_pkgconfig;
+      fi
   cache: []
 
 

CMakeLists.txt

@@ -52,14 +52,6 @@ set(GINKGO_CUDA_ARCHITECTURES "Auto" CACHE STRING
     "A list of target NVIDIA GPU architectures. See README.md for more detail.")
 # the details of fine/coarse grain memory and unsafe atomic are available https://docs.olcf.ornl.gov/systems/crusher_quick_start_guide.html#floating-point-fp-atomic-operations-and-coarse-fine-grained-memory-allocations
 option(GINKGO_HIP_AMD_UNSAFE_ATOMIC "Compiler uses unsafe floating point atomic (only for AMD GPU and ROCM >= 5). Default is ON because we use hipMalloc, which is always on coarse grain. Must turn off when allocating memory on fine grain" ON)
-set(GINKGO_HIP_COMPILER_FLAGS "" CACHE STRING
-    "Set the required HIP compiler flags. Current default is an empty string.")
-set(GINKGO_HIP_NVCC_COMPILER_FLAGS "" CACHE STRING
-    "Set the required HIP nvcc compiler flags. Current default is an empty string.")
-set(GINKGO_HIP_CLANG_COMPILER_FLAGS "" CACHE STRING
-    "Set the required HIP CLANG compiler flags. Current default is an empty string.")
-set(GINKGO_HIP_AMDGPU "" CACHE STRING
-    "The amdgpu_target(s) variable passed to hipcc. The default is none (auto).")
 option(GINKGO_SPLIT_TEMPLATE_INSTANTIATIONS "Split template instantiations for slow-to-compile files. This improves parallel build performance" ON)
 mark_as_advanced(GINKGO_SPLIT_TEMPLATE_INSTANTIATIONS)
 option(GINKGO_JACOBI_FULL_OPTIMIZATIONS "Use all the optimizations for the CUDA Jacobi algorithm" OFF)
@@ -144,12 +136,6 @@ if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
     set(CMAKE_BUILD_TYPE Release CACHE STRING "Choose the type of build." FORCE)
 endif()
 
-if(BUILD_SHARED_LIBS)
-    set(GINKGO_STATIC_OR_SHARED SHARED)
-else()
-    set(GINKGO_STATIC_OR_SHARED STATIC)
-endif()
-
 # Ensure we have a debug postfix
 if(NOT DEFINED CMAKE_DEBUG_POSTFIX)
     set(CMAKE_DEBUG_POSTFIX "d")
@@ -481,22 +467,26 @@ add_custom_target(test_exportbuild
     COMMAND ${GINKGO_TEST_EXPORTBUILD_CMD}
     COMMENT "Running a test on Ginkgo's exported build directory.")
 
-add_custom_target(test_pkgconfig
-    COMMAND ${CMAKE_COMMAND} -G${CMAKE_GENERATOR} ${TOOLSET}
-    -S${GINKGO_TEST_PKGCONFIG_SRC_DIR}
-    -B${GINKGO_TEST_PKGCONFIG_BIN_DIR}
-    -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER}
-    -DCMAKE_CUDA_COMPILER=${CMAKE_CUDA_COMPILER}
-    -DCMAKE_CXX_FLAGS=${CMAKE_CXX_FLAGS}
-    # `--config cfg` is ignored by single-configuration generator.
-    # `$<CONFIG>` is always be the same as `CMAKE_BUILD_TYPE` in
-    # single-configuration generator.
-    COMMAND ${CMAKE_COMMAND}
-    --build ${GINKGO_TEST_PKGCONFIG_BIN_DIR}
-    --config $<CONFIG>
-    COMMAND ${GINKGO_TEST_PKGCONFIG_CMD}
-    COMMENT "Running a test on Ginkgo's PkgConfig"
-    "This requires installing Ginkgo first")
+# static linking with pkg-config is not possible with HIP, since
+# some linker information cannot be expressed in pkg-config files
+if (BUILD_SHARED_LIBS OR NOT GINKGO_BUILD_HIP)
+    add_custom_target(test_pkgconfig
+        COMMAND ${CMAKE_COMMAND} -G${CMAKE_GENERATOR} ${TOOLSET}
+        -S${GINKGO_TEST_PKGCONFIG_SRC_DIR}
+        -B${GINKGO_TEST_PKGCONFIG_BIN_DIR}
+        -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER}
+        -DCMAKE_CUDA_COMPILER=${CMAKE_CUDA_COMPILER}
+        -DCMAKE_CXX_FLAGS=${CMAKE_CXX_FLAGS}
+        # `--config cfg` is ignored by single-configuration generator.
+        # `$<CONFIG>` is always be the same as `CMAKE_BUILD_TYPE` in
+        # single-configuration generator.
+        COMMAND ${CMAKE_COMMAND}
+        --build ${GINKGO_TEST_PKGCONFIG_BIN_DIR}
+        --config $<CONFIG>
+        COMMAND ${GINKGO_TEST_PKGCONFIG_CMD}
+        COMMENT "Running a test on Ginkgo's PkgConfig"
+        "This requires installing Ginkgo first")
+endif()
 
 
 # Setup CPack

INSTALL.md

@@ -51,8 +51,8 @@ Ginkgo adds the following additional switches to control what is being built:
 *   `-DGINKGO_BUILD_HIP={ON, OFF}` builds optimized HIP versions of the kernels
     (requires HIP), default is `ON` if an installation of HIP could be detected,
     `OFF` otherwise.
-*   `-DGINKGO_HIP_AMDGPU="gpuarch1;gpuarch2"` the amdgpu_target(s) variable
-    passed to hipcc for the `hcc` HIP backend. The default is none (auto).
+*   `-DCMAKE_HIP_ARCHITECTURES="gpuarch1;gpuarch2"` the AMDGPU targets to be passed to the compiler.
+    If empty, compiler chooses based on the available GPUs.
 *   `-DGINKGO_BUILD_HWLOC={ON, OFF}` builds Ginkgo with HWLOC. Default is `OFF`.
 *   `-DGINKGO_BUILD_DOC={ON, OFF}` creates an HTML version of Ginkgo's documentation
     from inline comments in the code. The default is `OFF`.
@@ -181,22 +181,13 @@ 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` 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:
+Ginkgo relies on CMake to decide which compiler to use for HIP.
+To choose `nvcc` instead of the default ROCm `clang++`, set the corresponding
+environment variable:
 ```bash
-export HIP_PLATFORM=nvcc # or nvidia for ROCM >= 4.1
+export HIPCXX=nvcc
 ```
-
-#### Setting platform specific compilation flags
-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.
+Note that this option is currently not being tested in our CI pipelines.
 
 
 ### Third party libraries and packages

README.md

@@ -63,6 +63,7 @@ The Ginkgo HIP module has the following __additional__ requirements:
     * _AMD_ backend (using the `clang` compiler)
     * _10.1 <= CUDA < 11_ backend
 * if the hipFFT package is available, it is used to implement the FFT LinOps.
+* _cmake 3.21+_
 
 The Ginkgo DPC++(SYCL) module has the following __additional__ requirements:
 

benchmark/CMakeLists.txt

@@ -25,14 +25,9 @@ endfunction()
 
 function(ginkgo_benchmark_hipsparse_linops type def)
     add_library(hipsparse_linops_${type} utils/hip_linops.hip.cpp)
+    set_source_files_properties(utils/hip_linops.hip.cpp PROPERTIES LANGUAGE HIP)
     target_compile_definitions(hipsparse_linops_${type} PUBLIC ${def})
-    EXECUTE_PROCESS(COMMAND ${HIP_PATH}/bin/hipconfig --cpp_config OUTPUT_VARIABLE HIP_CXX_FLAGS)
-    set_target_properties(hipsparse_linops_${type} PROPERTIES COMPILE_FLAGS ${HIP_CXX_FLAGS})
-    # use Thrust C++ device just for compilation, we don't use thrust::complex in the benchmarks
-    target_compile_definitions(hipsparse_linops_${type} PUBLIC -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CPP)
-    target_include_directories(hipsparse_linops_${type} SYSTEM PRIVATE
-        ${HSA_HEADER} ${HIP_INCLUDE_DIRS}
-        ${HIPBLAS_INCLUDE_DIRS} ${HIPSPARSE_INCLUDE_DIRS})
+    target_include_directories(hipsparse_linops_${type} SYSTEM PRIVATE ${HIPBLAS_INCLUDE_DIRS} ${HIPSPARSE_INCLUDE_DIRS})
     target_link_libraries(hipsparse_linops_${type} Ginkgo::ginkgo ${HIPSPARSE_LIBRARIES})
 endfunction()
 
@@ -126,10 +121,8 @@ if (GINKGO_BUILD_HIP)
     ginkgo_benchmark_hipsparse_linops(s GKO_BENCHMARK_USE_SINGLE_PRECISION)
     ginkgo_benchmark_hipsparse_linops(z GKO_BENCHMARK_USE_DOUBLE_COMPLEX_PRECISION)
     ginkgo_benchmark_hipsparse_linops(c GKO_BENCHMARK_USE_SINGLE_COMPLEX_PRECISION)
+    set_source_files_properties(utils/hip_timer.hip.cpp PROPERTIES LANGUAGE HIP)
     add_library(hip_timer utils/hip_timer.hip.cpp)
-    EXECUTE_PROCESS(COMMAND ${HIP_PATH}/bin/hipconfig --cpp_config OUTPUT_VARIABLE HIP_CXX_FLAGS)
-    set_target_properties(hip_timer PROPERTIES COMPILE_FLAGS ${HIP_CXX_FLAGS})
-    target_include_directories(hip_timer SYSTEM PRIVATE ${HSA_HEADER} ${HIP_INCLUDE_DIRS})
     target_link_libraries(hip_timer ginkgo)
 endif()
 

cmake/GinkgoConfig.cmake.in

@@ -63,16 +63,10 @@ set(GINKGO_JACOBI_FULL_OPTIMIZATIONS @GINKGO_JACOBI_FULL_OPTIMIZATIONS@)
 set(GINKGO_CUDA_ARCHITECTURES "@CMAKE_CUDA_ARCHITECTURES@")
 set(GINKGO_CUDA_HOST_COMPILER "@CMAKE_CUDA_HOST_COMPILER@")
 
-set(GINKGO_HIP_COMPILER_FLAGS "@GINKGO_HIP_COMPILER_FLAGS@")
-set(GINKGO_HIP_HCC_COMPILER_FLAGS "@GINKGO_HIP_HCC_COMPILER_FLAGS@")
-set(GINKGO_HIP_NVCC_COMPILER_FLAGS "@GINKGO_HIP_NVCC_COMPILER_FLAGS@")
-set(GINKGO_HIP_CLANG_COMPILER_FLAGS "@GINKGO_HIP_CLANG_COMPILER_FLAGS@")
-set(GINKGO_HIP_PLATFORM @GINKGO_HIP_PLATFORM@)
-set(GINKGO_HIP_PLATFORM_AMD_REGEX "@HIP_PLATFORM_AMD_REGEX@")
-set(GINKGO_HIP_PLATFORM_NVIDIA_REGEX "@HIP_PLATFORM_NVIDIA_REGEX@")
-set(GINKGO_HIP_AMDGPU "@GINKGO_HIP_AMDGPU@")
-set(GINKGO_HIP_VERSION @GINKGO_HIP_VERSION@)
-set(GINKGO_AMD_ARCH_FLAGS "@GINKGO_AMD_ARCH_FLAGS@")
+set(GINKGO_HIP_COMPILER_FLAGS "@CMAKE_HIP_COMPILER_FLAGS@")
+set(GINKGO_HIP_PLATFORM "@GINKGO_HIP_PLATFORM@")
+set(GINKGO_HIP_VERSION "@GINKGO_HIP_VERSION@")
+set(GINKGO_HIP_ARCHITECTURES "@CMAKE_HIP_ARCHITECTURES@")
 
 set(GINKGO_DPCPP_VERSION @GINKGO_DPCPP_VERSION@)
 set(GINKGO_DPCPP_MAJOR_VERSION @GINKGO_DPCPP_MAJOR_VERSION@)
@@ -175,7 +169,7 @@ if((NOT GINKGO_BUILD_SHARED_LIBS) AND GINKGO_BUILD_CUDA)
 endif()
 
 if((NOT GINKGO_BUILD_SHARED_LIBS) AND GINKGO_BUILD_HIP)
-    find_dependency(HIP)
+    enable_language(HIP)
     find_dependency(hipblas)
     find_dependency(hipfft)
     find_dependency(hiprand)
@@ -217,6 +211,9 @@ if(GINKGO_BUILD_CUDA)
     _ginkgo_check_compiler(CUDA)
     _ginkgo_check_compiler(CUDA_HOST)
 endif()
+if(GINKGO_BUILD_HIP)
+    _ginkgo_check_compiler(HIP)
+endif()
 
 include(${CMAKE_CURRENT_LIST_DIR}/GinkgoTargets.cmake)
 

cmake/autodetect_executors.cmake

@@ -31,7 +31,8 @@ if (NOT DEFINED GINKGO_BUILD_CUDA)
 endif()
 
 if (NOT DEFINED GINKGO_BUILD_HIP)
-    if(GINKGO_HIPCONFIG_PATH)
+    check_language(HIP)
+    if(CMAKE_HIP_COMPILER)
         message(STATUS "Enabling HIP executor")
         set(GINKGO_HAS_HIP ON)
     endif()

cmake/build_helpers.cmake

@@ -19,6 +19,10 @@ endfunction()
 
 function(ginkgo_compile_features name)
     target_compile_features("${name}" PUBLIC cxx_std_14)
+    # we set these properties regardless of the enabled backends,
+    # because unknown properties are ignored
+    set_target_properties("${name}" PROPERTIES HIP_STANDARD 14)
+    set_target_properties("${name}" PROPERTIES CUDA_STANDARD 14)
     if(GINKGO_WITH_CLANG_TIDY AND GINKGO_CLANG_TIDY_PATH)
         set_property(TARGET "${name}" PROPERTY CXX_CLANG_TIDY "${GINKGO_CLANG_TIDY_PATH};-checks=*")
     endif()
@@ -93,17 +97,12 @@ function(ginkgo_check_headers target defines)
         list(APPEND HIP_SOURCES "${HEADER_SOURCEFILE}")
     endforeach()
     if(HIP_SOURCES)
-        set_source_files_properties(${HIP_SOURCES} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT TRUE)
-        hip_add_library(${target}_headers_hip ${HIP_SOURCES}) # the compiler options get set by linking to ginkgo_hip
+        set_source_files_properties(${HIP_SOURCES} PROPERTIES LANGUAGE HIP)
+        add_library(${target}_headers_hip ${HIP_SOURCES}) # the compiler options get set by linking to ginkgo_hip
         target_link_libraries(${target}_headers_hip PRIVATE ${target} roc::hipblas roc::hipsparse hip::hiprand roc::rocrand)
         target_include_directories(${target}_headers_hip
             PRIVATE
-            "${CMAKE_CURRENT_SOURCE_DIR}"
-            "${GINKGO_HIP_THRUST_PATH}"
-            "${HIPBLAS_INCLUDE_DIRS}"
-            "${hiprand_INCLUDE_DIRS}"
-            "${HIPSPARSE_INCLUDE_DIRS}"
-            "${ROCPRIM_INCLUDE_DIRS}")
+            "${CMAKE_CURRENT_SOURCE_DIR}")
     endif()
 endfunction()
 

cmake/build_type_helpers.cmake

@@ -93,7 +93,7 @@ foreach(_LANG IN LISTS ENABLED_LANGUAGES ITEMS "HIP")
             set(${PROJECT_NAME}_${_LANG}_${_TYPE}_SUPPORTED FALSE)
         endif()
         if(${PROJECT_NAME}_${_LANG}_${_TYPE}_SUPPORTED)
-            if(_LANG STREQUAL "HIP" AND GINKGO_HIP_PLATFORM MATCHES "${HIP_PLATFORM_NVIDIA_REGEX}")
+            if(_LANG STREQUAL "HIP" AND GINKGO_HIP_PLATFORM_NVIDIA)
                 set(CMAKE_${_LANG}_FLAGS_${_TYPE}
                     ${${PROJECT_NAME}_NVCC_${_TYPE}_COMPILER_FLAGS}
                     CACHE STRING "Flags used by the ${_LANG} compiler during ${_TYPE} builds." FORCE