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 else()
 169     if(CMAKE_CXX_COMPILER MATCHES "dpcpp")
 170         # At least Intel dpcpp defaults to having SYCL enabled for all code. This leads to two problems:
 171         #
 172         # 1. Compiles take ~3x longer, since every file has to be compiled for multiple targets.
 173         # 2. We get a ton of warnings for the device-specific pass when the compiler sees our SIMD code.
 174         #
 175         # To avoid this, we attempt to find a flag to disable SYCL for non-SYCL files. Unfortunately,
 176         # when using gmx_find_flag_for_source() that includes calling check_cxx_compiler_flag(),
 177         # this in turn exposes a bug in dpcpp, where an object file compiles with -fno-sycl leads to
 178         # a failed link stage (when the same flag is not used). Since none of this is critical, we handle
 179         # it by merely checking if it works to compile a source fils with this flag, and choking if SYCL
 180         # is still enabled.
 181
 182         if(NOT CHECK_DISABLE_SYCL_CXX_FLAGS_QUIETLY)
 183             message(STATUS "Checking for flags to disable SYCL")
 184         endif()
 185
 186         gmx_check_source_compiles_with_flags(
 187             "int main() { return 0; }"
 188             "-fno-sycl"
 189             "CXX"
 190             DISABLE_SYCL_CXX_FLAGS_RESULT)
 191
 192         if(DISABLE_SYCL_CXX_FLAGS_RESULT)
 193             set(DISABLE_SYCL_CXX_FLAGS "-fno-sycl")
 194         endif()
 195         if(NOT CHECK_DISABLE_SYCL_CXX_FLAGS_QUIETLY)
 196             if(DISABLE_SYCL_CXX_FLAGS_RESULT)
 197                 message(STATUS "Checking for flags to disable SYCL - -fno-sycl")
 198             else()
 199                 message(WARNING "Cannot find flags to disable SYCL for non-SYCL hardware-specific C++ code. Expect many warnings, but they are likely benign.")
 200             endif()
 201             set(CHECK_DISABLE_SYCL_CXX_FLAGS_QUIETLY 1 CACHE INTERNAL "Keep quiet on future calls to detect no-SYCL flags" FORCE)
 202         endif()
 203     endif()
 204
 205     # Find the flags to enable (or re-enable) SYCL with Intel extensions. In case we turned it off above,
 206     # it's important that we check the combination of both flags, to make sure the second one re-enables SYCL.
 207     if(NOT CHECK_SYCL_CXX_FLAGS_QUIETLY)
 208         message(STATUS "Checking for flags to enable SYCL")
 209     endif()
 210     gmx_find_flag_for_source(SYCL_CXX_FLAGS_RESULT
 211         "#include <CL/sycl.hpp>
 212          namespace sycl = cl::sycl;
 213          int main(){
 214              sycl::queue q(sycl::default_selector{});
 215              return 0;
 216          }
 217          " "CXX" DISABLE_SYCL_CXX_FLAGS SYCL_CXX_FLAGS "-fsycl -fsycl-device-code-split=per_kernel")
 218
 219     if(NOT CHECK_SYCL_CXX_FLAGS_QUIETLY)
 220         if(SYCL_CXX_FLAGS_RESULT)
 221             message(STATUS "Checking for flags to enable SYCL - ${SYCL_CXX_FLAGS}")
 222         endif()
 223         set(CHECK_SYCL_CXX_FLAGS_QUIETLY 1 CACHE INTERNAL "Keep quiet on future calls to detect SYCL flags" FORCE)
 224     endif()
 225
 226     if(NOT SYCL_CXX_FLAGS_RESULT)
 227         message(FATAL_ERROR "Cannot compile with SYCL Intel compiler. Try a different compiler or disable SYCL.")
 228     endif()
 229
 230     # Add function wrapper similar to the one used by ComputeCPP and hipSYCL
 231     function(add_sycl_to_target)
 232         cmake_parse_arguments(
 233             PARSE_ARGV 0 # No positional arguments
 234             ARGS # Prefix for the resulting variables
 235             "" # No options
 236             "TARGET" # One-value keyword
 237             "SOURCES" # Multi-value keyword
 238         )
 239         set_source_files_properties(${ARGS_SOURCES} PROPERTIES COMPILE_FLAGS "${SYCL_CXX_FLAGS}")
 240         target_link_libraries(${ARGS_TARGET} PRIVATE ${SYCL_CXX_FLAGS})
 241     endfunction(add_sycl_to_target)
 242 endif()