Experimental MAS PCG solver

CMakeLists.txt

@@ -7,7 +7,7 @@ else()
 endif()
 
 # Check required CMake version
-set(REQUIRED_CMAKE_VERSION "3.18.0")
+set(REQUIRED_CMAKE_VERSION "3.24.0")
 if(POLYSOLVE_TOPLEVEL_PROJECT)
     cmake_minimum_required(VERSION ${REQUIRED_CMAKE_VERSION})
     SET(CMAKE_POLICY_VERSION_MINIMUM ${REQUIRED_CMAKE_VERSION})
@@ -75,6 +75,7 @@ endif()
 # Polysolve options
 option(POLYSOLVE_WITH_SANITIZERS    "Enable sanitizers in compilation targets"          OFF)
 option(POLYSOLVE_WITH_UNICODE       "Use Unicode in logging messages"                    ON)
+option(POLYSOLVE_WITH_CUDA          "Enable cuda support"                               OFF)
 
 # Polysolve options for enabling/disabling optional libraries
 option(POLYSOLVE_WITH_ACCELERATE    "Enable Apple Accelerate" ${POLYSOLVE_ON_APPLE_SILICON})
@@ -148,20 +149,79 @@ if(POLYSOLVE_WITH_MKL)
     set(CMAKE_MSVC_RUNTIME_LIBRARY "MultiThreadedDLL" CACHE STRING "Select the MSVC runtime library")
 endif()
 
+# Add empty library target. Setup later.
+add_library(polysolve)
+add_library(polysolve::polysolve ALIAS polysolve)
+add_library(polysolve_linear)
+add_library(polysolve::linear ALIAS polysolve_linear)
+
+################################################################################
+# CUDA
+################################################################################
+
+if (POLYSOLVE_WITH_CUDA)
+  # If CMAKE_CUDA_ARCHITECTURES was not specified, set it to native.
+  if(DEFINED CMAKE_CUDA_ARCHITECTURES)
+    message(STATUS "CMAKE_CUDA_ARCHITECTURES was specified, skipping auto-detection")
+  else()
+    message(STATUS "CMAKE_CUDA_ARCHITECTURES was not specified, set it to native")
+    set(CMAKE_CUDA_ARCHITECTURES "native")
+  endif()
+  message(STATUS "Targeting CUDA_ARCHITECTURES \"${CMAKE_CUDA_ARCHITECTURES}\"")
+
+  # Enable CUDA support
+  enable_language(CUDA)
+  set(CMAKE_CUDA_STANDARD 17)
+  set(CMAKE_CUDA_STANDARD_REQUIRED ON)
+
+  set_target_properties(polysolve PROPERTIES
+      CUDA_SEPARABLE_COMPILATION ON
+      CUDA_RESOLVE_DEVICE_SYMBOLS ON
+  )
+  set_target_properties(polysolve_linear PROPERTIES
+      CUDA_SEPARABLE_COMPILATION ON
+      CUDA_RESOLVE_DEVICE_SYMBOLS ON
+  )
+
+  # Enable __device__ lambda
+  target_compile_options(polysolve PRIVATE
+      $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>
+  )
+  target_compile_options(polysolve_linear PRIVATE
+      $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda>
+  )
+  # Relax device constexpr usage
+  target_compile_options(polysolve PRIVATE
+      $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>
+  )
+  target_compile_options(polysolve_linear PRIVATE
+      $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>
+  )
+
+  if(APPLE)
+      # We need to add the path to the driver (libcuda.dylib) as an rpath,
+      # so that the static cuda runtime can find it at runtime.
+      set_property(TARGET polysolve
+                  PROPERTY
+                  BUILD_RPATH ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
+  endif()
+endif()
+
 ################################################################################
 # PolySolve Library
 ################################################################################
 
-# Add an empty library and fill in the list of sources in `src/CMakeLists.txt`.
-add_library(polysolve_linear)
-add_library(polysolve::linear ALIAS polysolve_linear)
 add_subdirectory(src/polysolve)
 add_subdirectory(src/polysolve/linear)
 
-add_library(polysolve)
-add_library(polysolve::polysolve ALIAS polysolve)
 add_subdirectory(src/polysolve/nonlinear)
 
+# Ensure that changes in the CUDA linear library
+# trigger relinking of the aggregate polysolve
+# static library (and its CUDA device-link step).
+set_property(TARGET polysolve APPEND PROPERTY LINK_DEPENDS
+    "$<TARGET_FILE:polysolve_linear>")
+
 target_link_libraries(polysolve_linear PUBLIC polysolve_coverage_config)
 target_link_libraries(polysolve PUBLIC polysolve_coverage_config)
 
@@ -180,6 +240,11 @@ if(POLYSOLVE_LARGE_INDEX)
     target_compile_definitions(polysolve_linear PUBLIC POLYSOLVE_LARGE_INDEX)
 endif()
 
+if(POLYSOLVE_WITH_CUDA)
+  target_compile_definitions(polysolve PUBLIC POLYSOLVE_WITH_CUDA)
+  target_compile_definitions(polysolve_linear PUBLIC POLYSOLVE_WITH_CUDA)
+endif()
+
 target_compile_definitions(polysolve_linear PRIVATE POLYSOLVE_LINEAR_SPEC="${PROJECT_SOURCE_DIR}/linear-solver-spec.json")
 target_compile_definitions(polysolve PRIVATE POLYSOLVE_NON_LINEAR_SPEC="${PROJECT_SOURCE_DIR}/nonlinear-solver-spec.json")
 target_compile_definitions(polysolve_linear PUBLIC POLYSOLVE_JSON_SPEC_DIR="${PROJECT_SOURCE_DIR}")
@@ -189,6 +254,13 @@ target_compile_definitions(polysolve_linear PUBLIC POLYSOLVE_JSON_SPEC_DIR="${PR
 # Dependencies
 ################################################################################
 
+# Cuda runtime and cuSparse
+if (POLYSOLVE_WITH_CUDA)
+    find_package(CUDAToolkit REQUIRED)
+    target_link_libraries(polysolve PUBLIC CUDA::cudart CUDA::cusparse)
+    target_link_libraries(polysolve_linear PUBLIC CUDA::cudart CUDA::cusparse)
+endif()
+
 # ------
 # Linear
 # ------
@@ -229,8 +301,8 @@ if(POLYSOLVE_WITH_HYPRE)
     include(hypre)
     target_link_libraries(polysolve_linear PUBLIC HYPRE::HYPRE)
     target_compile_definitions(polysolve_linear PUBLIC POLYSOLVE_WITH_HYPRE)
-    if(HYPRE_WITH_MPI)
-        target_compile_definitions(polysolve_linear PUBLIC HYPRE_WITH_MPI)
+    if(HYPRE_ENABLE_MPI)
+        target_compile_definitions(polysolve_linear PUBLIC HYPRE_ENABLE_MPI)
     endif()
 endif()
 
@@ -304,6 +376,10 @@ endif()
 
 # cuSolver solvers
 if(POLYSOLVE_WITH_CUSOLVER)
+    if (NOT POLYSOLVE_WITH_CUDA)
+      message(FATAL_ERROR "Expect POLYSOLVE_WITH_CUDA set to ON when CUSOLVER is enabled")
+    endif()
+
     include(cusolverdn)
     if(TARGET CUDA::cusolver)
         target_link_libraries(polysolve_linear PUBLIC CUDA::cusolver)
@@ -331,6 +407,21 @@ if(POLYSOLVE_WITH_UNICODE)
     endif()
 endif()
 
+# CCCL
+if(POLYSOLVE_WITH_CUDA)
+  include(cccl)
+  target_link_libraries(polysolve_linear PUBLIC CCCL::CCCL)
+  target_link_libraries(polysolve PUBLIC CCCL::CCCL)
+endif()
+
+# KaMinPar
+# KaMinPar requires onetbb
+if(POLYSOLVE_WITH_CUDA)
+  include(kaminpar)
+  target_link_libraries(polysolve_linear PRIVATE KaMinPar::KaMinPar)
+endif()
+
+
 # ---------
 # Nonlinear
 # ---------

cmake/polysolve/polysolve_warnings.cmake

@@ -10,7 +10,6 @@ endif()
 set(POLYSOLVE_WARNING_FLAGS
 	-Wall
 	-Wextra
-	-pedantic
 
 	# -Wconversion
 	#-Wunsafe-loop-optimizations # broken with C++11 loops

cmake/recipes/amgcl.cmake

@@ -48,6 +48,18 @@ function(amgcl_import_target)
     )
 
     target_link_libraries(amgcl INTERFACE Boost::boost)
+
+    # AMGCL pollutes all downstream target with their aggressive diagostic flags,
+    # which causes cuda code to vomit warnings every line cause nvcc generates gcc specific intrinsics.
+    # Use hacky regex replacement to purge -Wpedantic, -Wall, and -Wextra.
+    get_target_property(AMGCL_OPTIONS amgcl INTERFACE_COMPILE_OPTIONS)
+    if(AMGCL_OPTIONS)
+        string(REGEX REPLACE ";?-Wpedantic" "" AMGCL_OPTIONS_CLEAN "${AMGCL_OPTIONS}")
+        string(REGEX REPLACE ";?-Wextra" "" AMGCL_OPTIONS_CLEAN "${AMGCL_OPTIONS_CLEAN}")
+        string(REGEX REPLACE ";?-Wall" "" AMGCL_OPTIONS_CLEAN "${AMGCL_OPTIONS_CLEAN}")
+        set_target_properties(amgcl PROPERTIES INTERFACE_COMPILE_OPTIONS "${AMGCL_OPTIONS_CLEAN}")
+    endif()
+
 endfunction()
 
 amgcl_import_target()

cmake/recipes/cccl.cmake

@@ -0,0 +1,15 @@
+# cccl (https://github.com/NVIDIA/cccl)
+# License: Apache
+
+if(TARGET CCCL::CCCL)
+    return()
+endif()
+
+message(STATUS "Third-party: creating target 'CCCL::CCCL'")
+
+include(CPM)
+CPMAddPackage(
+    NAME CCCL
+    GITHUB_REPOSITORY NVIDIA/cccl
+    VERSION 3.3.0
+)

cmake/recipes/cusolverdn.cmake

@@ -6,55 +6,11 @@ endif()
 
 message(STATUS "Third-party: creating targets 'CUDA::cusolver'")
 
-include(CheckLanguage)
-check_language(CUDA)
-if(CMAKE_CUDA_COMPILER)
-    enable_language(CUDA)
-
-    # We do not have a build recipe for this, so find it as a system installed library.
-    find_package(CUDAToolkit)
-    if(CUDAToolkit_FOUND)
-        set(CUDA_SEPARABLE_COMPILATION ON)
-        set(CUDA_PROPAGATE_HOST_FLAGS OFF)
-        set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -use_fast_math --expt-relaxed-constexpr  -gencode arch=compute_86,code=sm_86")
-
-        target_compile_options(polysolve PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:
-
-        --generate-line-info
-        --use_fast_math
-        --relocatable-device-code=true
-
-        --ptxas-options=-v
-        --maxrregcount=7
-        --compiler-options
-        -fPIC # https://stackoverflow.com/questions/5311515/gcc-fpic-option
-
-        >
-
-        )
-
-        target_link_libraries(polysolve PUBLIC CUDA::toolkit)
-        set_target_properties(polysolve PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
-        # Nvidia RTX8000 -> compute_75
-        # Nvidia V100 -> compute_70
-        # Nvidia 1080/1080Ti -> compute_61
-        # Nvidia 3080Ti -> compute_86
-        if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
-            set(CMAKE_CUDA_ARCHITECTURES 70 75 86)
-        endif()
-        set_target_properties(polysolve PROPERTIES CUDA_ARCHITECTURES "70;75;86")
-
-        if(APPLE)
-            # We need to add the path to the driver (libcuda.dylib) as an rpath,
-            # so that the static cuda runtime can find it at runtime.
-            set_property(TARGET polysolve
-                        PROPERTY
-                        BUILD_RPATH ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
-        endif()
-
-    else()
-        message(WARNING "cuSOLVER not found, solver will not be available.")
-    endif()
+# Use cuSolver bundled with cuda-toolkit.
+find_package(CUDAToolkit)
+if (CUDAToolkit_FIND)
+    target_link_libraries(polysolve PUBLIC CUDA::toolkit)
+    target_link_libraries(polysolve_linear PUBLIC CUDA::toolkit)
 else()
-    message(WARNING "CUDA not found, cuSOLVER will not be available.")
-endif()
+    message(WARNING "cuSOLVER not found, solver will not be available.")
+endif()

cmake/recipes/hypre.cmake

@@ -6,18 +6,28 @@ endif()
 
 message(STATUS "Third-party: creating target 'HYPRE::HYPRE'")
 
-set(HYPRE_WITH_MPI    OFF CACHE INTERNAL "" FORCE)
-set(HYPRE_PRINT_ERRORS  ON CACHE INTERNAL "" FORCE)
-set(HYPRE_BIGINT        ON CACHE INTERNAL "" FORCE)
-set(HYPRE_USING_FEI    OFF CACHE INTERNAL "" FORCE)
-set(HYPRE_USING_OPENMP OFF CACHE INTERNAL "" FORCE)
-set(HYPRE_SHARED       OFF CACHE INTERNAL "" FORCE)
+set(HYPRE_ENABLE_MPI           OFF CACHE INTERNAL "" FORCE)
+set(HYPRE_ENABLE_PRINT_ERRORS  ON  CACHE INTERNAL "" FORCE)
+set(HYPRE_ENABLE_BIGINT        ON  CACHE INTERNAL "" FORCE)
+set(HYPRE_ENABLE_FEI           OFF CACHE INTERNAL "" FORCE)
+set(HYPRE_ENABLE_OPENMP        OFF CACHE INTERNAL "" FORCE)
+set(HYPRE_SHARED               OFF CACHE INTERNAL "" FORCE)
+
+# HYPRE unconditionally defines an "uninstall" target, which conflicts with other buggy libraries
+# as modern cmake requires unique target name. This is a hacky workaround until upstream is fixed.
+macro(add_custom_target _target_name)
+  if("${_target_name}" STREQUAL "uninstall" AND TARGET uninstall)
+    # skip: HYPRE's uninstall target conflicts with an existing one
+  else()
+    _add_custom_target(${_target_name} ${ARGN})
+  endif()
+endmacro()
 
 include(CPM)
 CPMAddPackage(
     NAME hypre
     GITHUB_REPOSITORY hypre-space/hypre
-    GIT_TAG v2.28.0
+    GIT_TAG v3.1.0
     SOURCE_SUBDIR src
 )
 file(REMOVE "${hypre_SOURCE_DIR}/src/utilities/version")

cmake/recipes/kaminpar.cmake

@@ -0,0 +1,29 @@
+# KaMinPar (https://github.com/KaHIP/KaMinPar)
+# License: MIT
+
+if(TARGET KaMinPar::KaMinPar)
+    return()
+endif()
+
+message(STATUS "Third-party: creating target 'KaMinPar::KaMinPar'")
+
+include(onetbb)
+
+include(CPM)
+CPMAddPackage(
+    NAME KaMinPar
+    VERSION 3.7.3
+    GITHUB_REPOSITORY KaHIP/KaMinPar
+    GIT_TAG v3.7.3
+    OPTIONS
+        "KAMINPAR_BUILD_APPS OFF"
+        "KAMINPAR_BUILD_IO OFF"
+        "KAMINPAR_BUILD_WITH_SPARSEHASH OFF"
+        "KAMINPAR_BUILD_WITH_DEBUG_SYMBOLS OFF"
+        "KAMINPAR_64BIT_WEIGHTS ON"
+        "KAMINPAR_ENABLE_TIMERS OFF"
+        "INSTALL_KAMINPAR OFF"
+        # Below hack force kaminpar to use our vendored tbb
+        "KAMINPAR_DOWNLOAD_TBB ON"
+        "TBB_FOUND TRUE"
+)