cmake/gmxManageSYCL.cmake

   1 #
   2 # This file is part of the GROMACS molecular simulation package.
   3 #
   4 # Copyright (c) 2020,2021, by the GROMACS development team, led by
   5 # Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
   6 # and including many others, as listed in the AUTHORS file in the
   7 # top-level source directory and at http://www.gromacs.org.
   8 #
   9 # GROMACS is free software; you can redistribute it and/or
  10 # modify it under the terms of the GNU Lesser General Public License
  11 # as published by the Free Software Foundation; either version 2.1
  12 # of the License, or (at your option) any later version.
  13 #
  14 # GROMACS is distributed in the hope that it will be useful,
  15 # but WITHOUT ANY WARRANTY; without even the implied warranty of
  16 # MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  17 # Lesser General Public License for more details.
  18 #
  19 # You should have received a copy of the GNU Lesser General Public
  20 # License along with GROMACS; if not, see
  21 # http://www.gnu.org/licenses, or write to the Free Software Foundation,
  22 # Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301  USA.
  23 #
  24 # If you want to redistribute modifications to GROMACS, please
  25 # consider that scientific software is very special. Version
  26 # control is crucial - bugs must be traceable. We will be happy to
  27 # consider code for inclusion in the official distribution, but
  28 # derived work must not be called official GROMACS. Details are found
  29 # in the README & COPYING files - if they are missing, get the
  30 # official version at http://www.gromacs.org.
  31 #
  32 # To help us fund GROMACS development, we humbly ask that you cite
  33 # the research papers on the package. Check out http://www.gromacs.org.
  34
  35 # OpenCL required version: 1.2 or newer
  36 set(REQUIRED_SYCL_MIN_VERSION_MAJOR 1)
  37 set(REQUIRED_SYCL_MIN_VERSION_MINOR 2)
  38 set(REQUIRED_SYCL_MIN_VERSION ${REQUIRED_SYCL_MIN_VERSION_MAJOR}.${REQUIRED_SYCL_MIN_VERSION_MINOR})
  39
  40 set(GMX_GPU_SYCL ON)
  41
  42 # CMake issue tracking the efforts to make a universal upstream module:
  43 # https://gitlab.kitware.com/cmake/cmake/-/issues/21711
  44
  45 option(GMX_SYCL_HIPSYCL "Use hipSYCL instead of Intel/Clang for SYCL compilation" OFF)
  46
  47 if(GMX_DOUBLE)
  48     message(FATAL_ERROR "SYCL acceleration is not available in double precision")
  49 endif()
  50
  51 include(gmxFindFlagsForSource)
  52
  53 # Return all current CMake variables with name starting with "hipsycl" (case-insensitive).
  54 # Result is in the form of a list of flags ("-Dfoo=bar;-Dbaz=true").
  55 # Semicolons in values are escaped (needed for HIPSYCL_TARGETS).
  56 function(_getHipSyclCmakeFlags RETURN_VAR)
  57     get_cmake_property(_VARS VARIABLES)
  58     list (SORT _VARS)
  59     set(RESULT "")
  60     foreach (_VARNAME ${_VARS})
  61             string(TOLOWER "${_VARNAME}" _VARNAME_LOWER)
  62             if (${_VARNAME_LOWER} MATCHES "^hipsycl")
  63                 # Escape semicolon. The number of backslashes was determined empirically.
  64                 string(REPLACE ";" "\\\\\\;" _VARVALUE "${${_VARNAME}}")
  65                 list(APPEND
  66                   RESULT
  67                   -D${_VARNAME}=${_VARVALUE}
  68                 )
  69             endif()
  70     endforeach()
  71     set("${RETURN_VAR}" ${RESULT} PARENT_SCOPE)
  72 endfunction()
  73
  74 if(GMX_SYCL_HIPSYCL)
  75     set(HIPSYCL_CLANG "${CMAKE_CXX_COMPILER}")
  76     # -Wno-unknown-cuda-version because Clang-11 complains about CUDA 11.0-11.2, despite working fine with them.
  77     # -Wno-unknown-attributes because hipSYCL does not support reqd_sub_group_size (because it can only do some sub group sizes).
  78     set(HIPSYCL_SYCLCC_EXTRA_ARGS "-Wno-unknown-cuda-version -Wno-unknown-attributes")
  79
  80     # Must be called before find_package to capture all user-set CMake variables, but not those set automatically
  81     _getHipSyclCmakeFlags(_ALL_HIPSYCL_CMAKE_FLAGS)
  82
  83     find_package(hipsycl REQUIRED)
  84     # Ensure the try_compile projects below find the same hipsycl)
  85     list(APPEND _ALL_HIPSYCL_CMAKE_FLAGS -Dhipsycl_DIR=${hipsycl_DIR})
  86
  87     # If the user-set CMake variables change (e.g. because the user
  88     # changed HIPSYCL_TARGETS), then the try_compile tests below need
  89     # to be re-run. Set and use an internal cache variable to detect
  90     # the change and set a flag to rerun the tests.
  91     if (DEFINED GMX_ALL_HIPSYCL_CMAKE_FLAGS_COPY AND "${GMX_ALL_HIPSYCL_CMAKE_FLAGS_COPY}" STREQUAL "${_ALL_HIPSYCL_CMAKE_FLAGS}")
  92         set(_rerun_hipsycl_try_compile_tests FALSE)
  93     else()
  94         # The new value should over-write the previous copy
  95         set(GMX_ALL_HIPSYCL_CMAKE_FLAGS_COPY ${_ALL_HIPSYCL_CMAKE_FLAGS} CACHE INTERNAL "Store the list of CMake variables needed for hipSYCL compilation test projects")
  96         set(_rerun_hipsycl_try_compile_tests TRUE)
  97     endif()
  98
  99     # Does the hipSYCL compiler work at all for the given targets?
 100     if (NOT DEFINED GMX_HIPSYCL_COMPILATION_WORKS OR _rerun_hipsycl_try_compile_tests)
 101         message(STATUS "Checking for valid hipSYCL compiler")
 102         try_compile(GMX_HIPSYCL_COMPILATION_WORKS "${CMAKE_BINARY_DIR}/CMakeTmpHipSyclTest" "${CMAKE_SOURCE_DIR}/cmake/HipSyclTest/" "HipSyclTest"
 103           OUTPUT_VARIABLE _HIPSYCL_COMPILATION_OUTPUT
 104           CMAKE_FLAGS
 105             ${_ALL_HIPSYCL_CMAKE_FLAGS})
 106         file(REMOVE_RECURSE "${CMAKE_BINARY_DIR}/CMakeTmpHipSyclTest")
 107         if(GMX_HIPSYCL_COMPILATION_WORKS)
 108             message(STATUS "Checking for valid hipSYCL compiler - Success")
 109         endif()
 110     endif()
 111     if (NOT GMX_HIPSYCL_COMPILATION_WORKS)
 112         message(FATAL_ERROR "hipSYCL compiler not working:\n${_HIPSYCL_COMPILATION_OUTPUT}")
 113     endif()
 114
 115     # Does hipSYCL compilation target CUDA devices?
 116     if(NOT DEFINED GMX_HIPSYCL_HAVE_CUDA_TARGET OR _rerun_hipsycl_try_compile_tests)
 117         message(STATUS "Checking for hipSYCL CUDA target")
 118         try_compile(GMX_HIPSYCL_HAVE_CUDA_TARGET "${CMAKE_BINARY_DIR}/CMakeTmpHipSyclTest" "${CMAKE_SOURCE_DIR}/cmake/HipSyclTest/" "HipSyclTest"
 119           CMAKE_FLAGS
 120             -DCHECK_CUDA_TARGET=ON
 121             ${_ALL_HIPSYCL_CMAKE_FLAGS})
 122         file(REMOVE_RECURSE "${CMAKE_BINARY_DIR}/CMakeTmpHipSyclTest")
 123         if(GMX_HIPSYCL_HAVE_CUDA_TARGET)
 124             message(STATUS "Checking for hipSYCL CUDA target - Success")
 125         else()
 126             message(STATUS "Checking for hipSYCL CUDA target - Failed")
 127         endif()
 128     endif()
 129
 130     # Does hipSYCL compilation target HIP devices?
 131     if(NOT DEFINED GMX_HIPSYCL_HAVE_HIP_TARGET OR _rerun_hipsycl_try_compile_tests)
 132         message(STATUS "Checking for hipSYCL HIP target")
 133         try_compile(GMX_HIPSYCL_HAVE_HIP_TARGET "${CMAKE_BINARY_DIR}/CMakeTmpHipSyclTest" "${CMAKE_SOURCE_DIR}/cmake/HipSyclTest/" "HipSyclTest"
 134           CMAKE_FLAGS
 135             -DCHECK_HIP_TARGET=ON
 136             ${_ALL_HIPSYCL_CMAKE_FLAGS})
 137         file(REMOVE_RECURSE "${CMAKE_BINARY_DIR}/CMakeTmpHipSyclTest")
 138         if(GMX_HIPSYCL_HAVE_HIP_TARGET)
 139             message(STATUS "Checking for hipSYCL HIP target - Success")
 140         else()
 141             message(STATUS "Checking for hipSYCL HIP target - Failed")
 142         endif()
 143     endif()
 144
 145     # Does hipSYCL compilation target Intel Level0 devices?
 146     if(NOT DEFINED GMX_HIPSYCL_HAVE_LEVELZERO_TARGET OR _rerun_hipsycl_try_compile_tests)
 147         message(STATUS "Checking for hipSYCL LevelZero target")
 148         try_compile(GMX_HIPSYCL_HAVE_LEVELZERO_TARGET "${CMAKE_BINARY_DIR}/CMakeTmpHipSyclTest" "${CMAKE_SOURCE_DIR}/cmake/HipSyclTest/" "HipSyclTest"
 149           CMAKE_FLAGS
 150             -DCHECK_LEVELZERO_TARGET=ON
 151             ${_ALL_HIPSYCL_CMAKE_FLAGS})
 152         file(REMOVE_RECURSE "${CMAKE_BINARY_DIR}/CMakeTmpHipSyclTest")
 153         if(GMX_HIPSYCL_HAVE_LEVELZERO_TARGET)
 154             message(STATUS "Checking for hipSYCL LevelZero target - Success")
 155             message(WARNING "GROMACS does not support LevelZero backend of hipSYCL")
 156         else()
 157             message(STATUS "Checking for hipSYCL LevelZero target - Failed")
 158         endif()
 159     endif()
 160
 161     if(NOT GMX_HIPSYCL_HAVE_CUDA_TARGET AND NOT GMX_HIPSYCL_HAVE_HIP_TARGET)
 162         message(WARNING "hipSYCL has no GPU targets set! Please, specify target hardware with -DHIPSYCL_TARGETS CMake option")
 163     endif()
 164     if(GMX_HIPSYCL_HAVE_CUDA_TARGET AND GMX_HIPSYCL_HAVE_HIP_TARGET)
 165         message(FATAL_ERROR "hipSYCL cannot have both CUDA and HIP targets active! This would require explicit multipass mode which both decreases performance on NVIDIA devices and has been removed in clang 12. Compile only for either CUDA or HIP targets.")
 166     endif()
 167     unset(_rerun_hipsycl_try_compile_tests)
 168
 169     # Find a suitable rocFFT when hipSYCL is targeting AMD devices
 170     if (GMX_HIPSYCL_HAVE_HIP_TARGET)
 171         # For consistency, we prefer to find rocFFT as part of the
 172         # default ROCm distribution that supports the version of
 173         # hipSYCL that is being used. Other installations of rocFFT
 174         # might work, but could lead to problems that are hard to
 175         # trace.
 176         #
 177         # The hipSYCL find package sets HIPSYCL_SYCLCC which we can
 178         # use to find the JSON configuration file that points to the
 179         # default ROCm installation used by hipSYCL, which can be used
 180         # to find rocFFT.
 181         #
 182         # If this is unavailable or does not work, the user will need to
 183         # set CMAKE_PREFIX_PATH so CMake is able to find the dependencies
 184         # of rocFFT (namely hip, AMDDeviceLibs, amd_comgr, hsa-runtime64,
 185         # ROCclr).
 186         if (HIPSYCL_SYCLCC)
 187             get_filename_component(HIPSYCL_SYCLCC_DIR ${HIPSYCL_SYCLCC} DIRECTORY)
 188             find_file(HIPSYCL_SYCLCC_JSON syclcc.json
 189                 HINTS ${HIPSYCL_SYCLCC_DIR}/../etc/hipSYCL
 190                 DOC "location of hipSYCL JSON configuration file"
 191                 )
 192             if (HIPSYCL_SYCLCC_JSON)
 193                 if(NOT HIPSYCL_SYCLCC_ROCM_PATH)
 194                     file(READ ${HIPSYCL_SYCLCC_JSON} HIPSYCL_SYCLCC_JSON_CONTENTS)
 195                     if (CMAKE_VERSION VERSION_LESS 3.19)
 196                         # We want the value encoded by the line
 197                         # "default-rocm-path" : "/opt/rocm",
 198                         # so we use regular expressions to remove everything before
 199                         # and after the relevant quotation marks.
 200                         #
 201                         # Remove this when GROMACS requires CMake 3.19 or higher, as the
 202                         # proper JSON parsing below is more robust.
 203                         string(REGEX REPLACE ".*\"default-rocm-path\" *: * \"" "" HIPSYCL_SYCLCC_ROCM_PATH_VALUE ${HIPSYCL_SYCLCC_JSON_CONTENTS})
 204                         string(REGEX REPLACE "\",.*" "" HIPSYCL_SYCLCC_ROCM_PATH_VALUE ${HIPSYCL_SYCLCC_ROCM_PATH_VALUE})
 205                     else()
 206                         string(JSON HIPSYCL_SYCLCC_ROCM_PATH_VALUE GET ${HIPSYCL_SYCLCC_JSON_CONTENTS} "default-rocm-path")
 207                     endif()
 208                     set(HIPSYCL_SYCLCC_ROCM_PATH ${HIPSYCL_SYCLCC_ROCM_PATH_VALUE} CACHE FILEPATH "The default ROCm used by syclcc from hipSYCL")
 209                 endif()
 210
 211                 if(HIPSYCL_SYCLCC_ROCM_PATH)
 212                     # Teach the rocFFT find package how to find the necessary components
 213                     # from the ROCm distribution used by hipSYCL.
 214                     set(hip_DIR ${HIPSYCL_SYCLCC_ROCM_PATH}/hip/lib/cmake/hip)
 215                     set(AMDDeviceLibs_DIR ${HIPSYCL_SYCLCC_ROCM_PATH}/lib/cmake/AMDDeviceLibs)
 216                     set(amd_comgr_DIR ${HIPSYCL_SYCLCC_ROCM_PATH}/lib/cmake/amd_comgr)
 217                     set(hsa-runtime64_DIR ${HIPSYCL_SYCLCC_ROCM_PATH}/lib/cmake/hsa-runtime64)
 218                     set(ROCclr_DIR ${HIPSYCL_SYCLCC_ROCM_PATH}/rocclr/lib/cmake/rocclr)
 219                     set(rocfft_DIR ${HIPSYCL_SYCLCC_ROCM_PATH}/rocfft/lib/cmake/rocfft)
 220                 endif()
 221             endif()
 222         endif()
 223
 224         # Find rocFFT, either from the ROCm used by hipSYCL, or as otherwise found on the system
 225         find_package(rocfft ${FIND_ROCFFT_QUIETLY} CONFIG HINTS ${HIPSYCL_SYCLCC_ROCM_PATH} PATHS /opt/rocm)
 226         if (NOT rocfft_FOUND)
 227             message(FATAL_ERROR "rocFFT is required for the hipSYCL build, but was not found")
 228         endif()
 229         set(FIND_ROCFFT_QUIETLY "QUIET")
 230     endif()
 231 else()
 232     if(CMAKE_CXX_COMPILER MATCHES "dpcpp")
 233         # At least Intel dpcpp defaults to having SYCL enabled for all code. This leads to two problems:
 234         #
 235         # 1. Compiles take ~3x longer, since every file has to be compiled for multiple targets.
 236         # 2. We get a ton of warnings for the device-specific pass when the compiler sees our SIMD code.
 237         #
 238         # To avoid this, we attempt to find a flag to disable SYCL for non-SYCL files. Unfortunately,
 239         # when using gmx_find_flag_for_source() that includes calling check_cxx_compiler_flag(),
 240         # this in turn exposes a bug in dpcpp, where an object file compiles with -fno-sycl leads to
 241         # a failed link stage (when the same flag is not used). Since none of this is critical, we handle
 242         # it by merely checking if it works to compile a source fils with this flag, and choking if SYCL
 243         # is still enabled.
 244
 245         if(NOT CHECK_DISABLE_SYCL_CXX_FLAGS_QUIETLY)
 246             message(STATUS "Checking for flags to disable SYCL")
 247         endif()
 248
 249         gmx_check_source_compiles_with_flags(
 250             "int main() { return 0; }"
 251             "-fno-sycl"
 252             "CXX"
 253             DISABLE_SYCL_CXX_FLAGS_RESULT)
 254
 255         if(DISABLE_SYCL_CXX_FLAGS_RESULT)
 256             set(DISABLE_SYCL_CXX_FLAGS "-fno-sycl")
 257         endif()
 258         if(NOT CHECK_DISABLE_SYCL_CXX_FLAGS_QUIETLY)
 259             if(DISABLE_SYCL_CXX_FLAGS_RESULT)
 260                 message(STATUS "Checking for flags to disable SYCL - -fno-sycl")
 261             else()
 262                 message(WARNING "Cannot find flags to disable SYCL for non-SYCL hardware-specific C++ code. Expect many warnings, but they are likely benign.")
 263             endif()
 264             set(CHECK_DISABLE_SYCL_CXX_FLAGS_QUIETLY 1 CACHE INTERNAL "Keep quiet on future calls to detect no-SYCL flags" FORCE)
 265         endif()
 266     endif()
 267
 268     # Find the flags to enable (or re-enable) SYCL with Intel extensions. In case we turned it off above,
 269     # it's important that we check the combination of both flags, to make sure the second one re-enables SYCL.
 270     if(NOT CHECK_SYCL_CXX_FLAGS_QUIETLY)
 271         message(STATUS "Checking for flags to enable SYCL")
 272     endif()
 273     gmx_find_flag_for_source(SYCL_CXX_FLAGS_RESULT
 274         "#include <CL/sycl.hpp>
 275          namespace sycl = cl::sycl;
 276          int main(){
 277              sycl::queue q(sycl::default_selector{});
 278              return 0;
 279          }
 280          " "CXX" DISABLE_SYCL_CXX_FLAGS SYCL_CXX_FLAGS "-fsycl -fsycl-device-code-split=per_kernel")
 281
 282     if(NOT CHECK_SYCL_CXX_FLAGS_QUIETLY)
 283         if(SYCL_CXX_FLAGS_RESULT)
 284             message(STATUS "Checking for flags to enable SYCL - ${SYCL_CXX_FLAGS}")
 285         endif()
 286         set(CHECK_SYCL_CXX_FLAGS_QUIETLY 1 CACHE INTERNAL "Keep quiet on future calls to detect SYCL flags" FORCE)
 287     endif()
 288
 289     if(NOT SYCL_CXX_FLAGS_RESULT)
 290         message(FATAL_ERROR "Cannot compile with SYCL Intel compiler. Try a different compiler or disable SYCL.")
 291     endif()
 292
 293     include(gmxManageFFTLibraries)
 294     if(NOT GMX_FFT_MKL)
 295         message(WARNING "Building SYCL version with ${GMX_FFT_LIBRARY} instead of MKL. GPU FFT is disabled!")
 296     endif()
 297
 298     # Add function wrapper similar to the one used by ComputeCPP and hipSYCL
 299     function(add_sycl_to_target)
 300         cmake_parse_arguments(
 301             PARSE_ARGV 0 # No positional arguments
 302             ARGS # Prefix for the resulting variables
 303             "" # No options
 304             "TARGET" # One-value keyword
 305             "SOURCES" # Multi-value keyword
 306         )
 307         set_source_files_properties(${ARGS_SOURCES} PROPERTIES COMPILE_FLAGS "${SYCL_CXX_FLAGS}")
 308         target_link_libraries(${ARGS_TARGET} PRIVATE ${SYCL_CXX_FLAGS})
 309     endfunction(add_sycl_to_target)
 310 endif()