339 lines · plain
1##===----------------------------------------------------------------------===##2#3# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4# See https://llvm.org/LICENSE.txt for license information.5# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6#7##===----------------------------------------------------------------------===##8#9# Find OpenMP Target offloading Support for various compilers.10#11##===----------------------------------------------------------------------===##12 13#[========================================================================[.rst:14FindOpenMPTarget15----------------16 17Finds OpenMP Target Offloading Support.18 19This module can be used to detect OpenMP target offloading support in a20compiler. If the compiler support OpenMP Offloading to a specified target, the21flags required to compile offloading code to that target are output for each22target.23 24This module will automatically include OpenMP support if it was not loaded25already. It does not need to be included separately to get full OpenMP support.26 27Variables28^^^^^^^^^29 30The module exposes the components ``NVPTX`` and ``AMDGPU``. Each of these31controls the various offloading targets to search OpenMP target offloasing32support for.33 34Depending on the enabled components the following variables will be set:35 36``OpenMPTarget_FOUND``37 Variable indicating that OpenMP target offloading flags for all requested38 targets have been found.39 40This module will set the following variables per language in your41project, where ``<device>`` is one of NVPTX or AMDGPU42 43``OpenMPTarget_<device>_FOUND``44 Variable indicating if OpenMP support for the ``<device>`` was detected.45``OpenMPTarget_<device>_FLAGS``46 OpenMP compiler flags for offloading to ``<device>``, separated by spaces.47 48For linking with OpenMP code written in ``<device>``, the following49variables are provided:50 51``OpenMPTarget_<device>_LIBRARIES``52 A list of libraries needed to link with OpenMP code written in ``<lang>``.53 54Additionally, the module provides :prop_tgt:`IMPORTED` targets:55 56``OpenMPTarget::OpenMPTarget_<device>``57 Target for using OpenMP offloading to ``<device>``.58 59If the specific architecture of the target is needed, it can be manually60specified by setting a variable to the desired architecture. Variables can also61be used to override the standard flag searching for a given compiler.62 63``OpenMPTarget_<device>_ARCH``64 Sets the architecture of ``<device>`` to compile for. Such as `sm_70` for NVPTX65 or `gfx908` for AMDGPU. 66 67``OpenMPTarget_<device>_DEVICE``68 Sets the name of the device to offload to.69 70``OpenMPTarget_<device>_FLAGS``71 Sets the compiler flags for offloading to ``<device>``.72 73#]========================================================================]74 75# TODO: Support Fortran76# TODO: Support multiple offloading targets by setting the "OpenMPTarget" target77# to include flags for all components loaded78# TODO: Configure target architecture without a variable (component NVPTX_SM_70)79# TODO: Test more compilers80 81cmake_policy(PUSH)82cmake_policy(VERSION 3.20.0)83 84find_package(OpenMP ${OpenMPTarget_FIND_VERSION} REQUIRED)85 86# Find the offloading flags for each compiler.87function(_OPENMP_TARGET_DEVICE_FLAG_CANDIDATES LANG DEVICE)88 if(NOT OpenMPTarget_${LANG}_FLAGS)89 unset(OpenMPTarget_FLAG_CANDIDATES)90 91 set(OMPTarget_FLAGS_Clang "-fopenmp-targets=${DEVICE}")92 set(OMPTarget_FLAGS_GNU "-foffload=${DEVICE}=\"-lm -latomic\"")93 set(OMPTarget_FLAGS_XL "-qoffload")94 set(OMPTarget_FLAGS_PGI "-mp=${DEVICE}")95 set(OMPTarget_FLAGS_NVHPC "-mp=${DEVICE}")96 97 if(DEFINED OMPTarget_FLAGS_${CMAKE_${LANG}_COMPILER_ID})98 set(OpenMPTarget_FLAG_CANDIDATES "${OMPTarget_FLAGS_${CMAKE_${LANG}_COMPILER_ID}}")99 endif()100 101 set(OpenMPTarget_${LANG}_FLAG_CANDIDATES "${OpenMPTarget_FLAG_CANDIDATES}" PARENT_SCOPE)102 else()103 set(OpenMPTarget_${LANG}_FLAG_CANDIDATES "${OpenMPTarget_${LANG}_FLAGS}" PARENT_SCOPE)104 endif()105endfunction()106 107# Get the coded name of the device for each compiler.108function(_OPENMP_TARGET_DEVICE_CANDIDATES LANG DEVICE)109 if (NOT OpenMPTarget_${DEVICE}_DEVICE)110 unset(OpenMPTarget_DEVICE_CANDIDATES)111 112 # Check each supported device.113 if("${DEVICE}" STREQUAL "NVPTX")114 if ("${CMAKE_SIZEOF_VOID_P}" STREQUAL "4")115 set(OMPTarget_DEVICE_Clang "nvptx32-nvidia-cuda")116 else()117 set(OMPTarget_DEVICE_Clang "nvptx64-nvidia-cuda")118 endif()119 set(OMPTarget_DEVICE_GNU "nvptx-none")120 set(OMPTarget_DEVICE_XL "")121 set(OMPTarget_DEVICE_PGI "gpu")122 set(OMPTarget_DEVICE_NVHPC "gpu")123 124 if(DEFINED OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID})125 set(OpenMPTarget_DEVICE_CANDIDATES "${OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID}}")126 endif()127 elseif("${DEVICE}" STREQUAL "AMDGPU")128 set(OMPTarget_DEVICE_Clang "amdgcn-amd-amdhsa")129 set(OMPTarget_DEVICE_GNU "hsa")130 131 if(DEFINED OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID})132 set(OpenMPTarget_DEVICE_CANDIDATES "${OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID}}")133 endif()134 endif()135 set(OpenMPTarget_${LANG}_DEVICE_CANDIDATES "${OpenMPTarget_DEVICE_CANDIDATES}" PARENT_SCOPE)136 else()137 set(OpenMPTarget_${LANG}_DEVICE_CANDIDATES "${OpenMPTarget_${LANG}_DEVICE}" PARENT_SCOPE)138 endif()139endfunction()140 141# Get flags for setting the device's architecture for each compiler.142function(_OPENMP_TARGET_DEVICE_ARCH_CANDIDATES LANG DEVICE DEVICE_FLAG)143 if(OpenMPTarget_${DEVICE}_ARCH)144 # Only Clang supports selecting the architecture for now.145 set(OMPTarget_ARCH_Clang "-Xopenmp-target=${DEVICE_FLAG} -march=${OpenMPTarget_${DEVICE}_ARCH}")146 147 if(DEFINED OMPTarget_ARCH_${CMAKE_${LANG}_COMPILER_ID})148 set(OpenMPTarget_DEVICE_ARCH_CANDIDATES "${OMPTarget_ARCH_${CMAKE_${LANG}_COMPILER_ID}}")149 endif()150 set(OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES "${OpenMPTarget_DEVICE_ARCH_CANDIDATES}" PARENT_SCOPE)151 else()152 set(OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES "" PARENT_SCOPE)153 endif()154endfunction()155 156set(OpenMPTarget_C_CXX_TEST_SOURCE157"#include <omp.h>158int main(void) {159 int isHost;160#pragma omp target map(from: isHost)161 { isHost = omp_is_initial_device(); }162 return isHost;163}")164 165function(_OPENMP_TARGET_WRITE_SOURCE_FILE LANG SRC_FILE_CONTENT_VAR SRC_FILE_NAME SRC_FILE_FULLPATH)166 set(WORK_DIR ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/FindOpenMPTarget)167 if("${LANG}" STREQUAL "C")168 set(SRC_FILE "${WORK_DIR}/${SRC_FILE_NAME}.c")169 file(WRITE "${SRC_FILE}" "${OpenMPTarget_C_CXX_${SRC_FILE_CONTENT_VAR}}")170 elseif("${LANG}" STREQUAL "CXX")171 set(SRC_FILE "${WORK_DIR}/${SRC_FILE_NAME}.cpp")172 file(WRITE "${SRC_FILE}" "${OpenMPTarget_C_CXX_${SRC_FILE_CONTENT_VAR}}")173 endif()174 set(${SRC_FILE_FULLPATH} "${SRC_FILE}" PARENT_SCOPE)175endfunction()176 177# Get the candidate flags and try to compile a test application. If it compiles178# and all the flags are found, we assume the compiler supports offloading.179function(_OPENMP_TARGET_DEVICE_GET_FLAGS LANG DEVICE OPENMP_FLAG_VAR OPENMP_LIB_VAR OPENMP_DEVICE_VAR OPENMP_ARCH_VAR)180 _OPENMP_TARGET_DEVICE_CANDIDATES(${LANG} ${DEVICE})181 _OPENMP_TARGET_DEVICE_FLAG_CANDIDATES(${LANG} "${OpenMPTarget_${LANG}_DEVICE_CANDIDATES}")182 _OPENMP_TARGET_DEVICE_ARCH_CANDIDATES(${LANG} ${DEVICE} "${OpenMPTarget_${LANG}_DEVICE_CANDIDATES}")183 _OPENMP_TARGET_WRITE_SOURCE_FILE("${LANG}" "TEST_SOURCE" OpenMPTargetTryFlag _OPENMP_TEST_SRC)184 185 # Try to compile a test application with the found flags.186 try_compile(OpenMPTarget_COMPILE_RESULT ${CMAKE_BINARY_DIR} ${_OPENMP_TEST_SRC}187 CMAKE_FLAGS "-DCOMPILE_DEFINITIONS:STRING=${OpenMP_${LANG}_FLAGS} ${OpenMPTarget_${LANG}_FLAG_CANDIDATES} ${OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES}"188 "-DINCLUDE_DIRECTORIES:STRING=${OpenMP_${LANG}_INCLUDE_DIR}"189 LINK_LIBRARIES ${CMAKE_${LANG}_VERBOSE_FLAG}190 OUTPUT_VARIABLE OpenMP_TRY_COMPILE_OUTPUT191 )192 193 file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log194 "Detecting OpenMP ${CMAKE_${LANG}_COMPILER_ID} ${DEVICE} target support with the following Flags:195 ${OpenMP_${LANG}_FLAGS} ${OpenMPTarget_${LANG}_FLAG_CANDIDATES} ${OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES}196 With the following output:\n ${OpenMP_TRY_COMPILE_OUTPUT}\n")197 198 # If compilation was successful and the device was found set the return variables.199 if (OpenMPTarget_COMPILE_RESULT AND DEFINED OpenMPTarget_${LANG}_DEVICE_CANDIDATES)200 file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log201 "Compilation successful, adding flags for ${DEVICE}.\n\n")202 203 # Clang has a separate library for target offloading.204 if(CMAKE_${LANG}_COMPILER_ID STREQUAL "Clang")205 find_library(OpenMPTarget_libomptarget_LIBRARY206 NAMES omptarget207 HINTS ${CMAKE_${LANG}_IMPLICIT_LINK_DIRECTORIES}208 )209 mark_as_advanced(OpenMPTarget_libomptarget_LIBRARY)210 set("${OPENMP_LIB_VAR}" "${OpenMPTarget_libomptarget_LIBRARY}" PARENT_SCOPE)211 else()212 unset("${OPENMP_LIB_VAR}" PARENT_SCOPE)213 endif()214 set("${OPENMP_DEVICE_VAR}" "${OpenMPTarget_${LANG}_DEVICE_CANDIDATES}" PARENT_SCOPE)215 set("${OPENMP_FLAG_VAR}" "${OpenMPTarget_${LANG}_FLAG_CANDIDATES}" PARENT_SCOPE)216 set("${OPENMP_ARCH_VAR}" "${OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES}" PARENT_SCOPE)217 else()218 unset("${OPENMP_DEVICE_VAR}" PARENT_SCOPE)219 unset("${OPENMP_FLAG_VAR}" PARENT_SCOPE)220 unset("${OPENMP_ARCH_VAR}" PARENT_SCOPE)221 endif()222endfunction()223 224# Load the compiler support for each device.225foreach(LANG IN ITEMS C CXX)226 # Cache the version in case CMake doesn't load the OpenMP package this time227 set(OpenMP_${LANG}_VERSION ${OpenMP_${LANG}_VERSION}228 CACHE STRING "OpenMP Version" FORCE)229 mark_as_advanced(OpenMP_${LANG}_VERSION)230 foreach(DEVICE IN ITEMS NVPTX AMDGPU)231 if(CMAKE_${LANG}_COMPILER_LOADED)232 if(NOT DEFINED OpenMPTarget_${LANG}_FLAGS OR NOT DEFINED OpenMPTarget_${LANG}_DEVICE)233 _OPENMP_TARGET_DEVICE_GET_FLAGS(${LANG} ${DEVICE}234 OpenMPTarget_${DEVICE}_FLAGS_WORK235 OpenMPTarget_${DEVICE}_LIBS_WORK236 OpenMPTarget_${DEVICE}_DEVICE_WORK237 OpenMPTarget_${DEVICE}_ARCHS_WORK)238 239 separate_arguments(_OpenMPTarget_${DEVICE}_FLAGS NATIVE_COMMAND "${OpenMPTarget_${DEVICE}_FLAGS_WORK}")240 separate_arguments(_OpenMPTarget_${DEVICE}_ARCHS NATIVE_COMMAND "${OpenMPTarget_${DEVICE}_ARCHS_WORK}")241 set(OpenMPTarget_${DEVICE}_FLAGS ${_OpenMPTarget_${DEVICE}_FLAGS}242 CACHE STRING "${DEVICE} target compile flags for OpenMP target offloading" FORCE)243 set(OpenMPTarget_${DEVICE}_ARCH ${_OpenMPTarget_${DEVICE}_ARCHS}244 CACHE STRING "${DEVICE} target architecture flags for OpenMP target offloading" FORCE)245 set(OpenMPTarget_${DEVICE}_LIBRARIES ${OpenMPTarget_${DEVICE}_LIBS_WORK}246 CACHE STRING "${DEVICE} target libraries for OpenMP target offloading" FORCE)247 mark_as_advanced(OpenMPTarget_${DEVICE}_FLAGS OpenMPTarget_${DEVICE}_ARCH OpenMPTarget_${DEVICE}_LIBRARIES)248 endif()249 endif()250 endforeach()251endforeach()252 253if(OpenMPTarget_FIND_COMPONENTS)254 set(OpenMPTarget_FINDLIST ${OpenMPTarget_FIND_COMPONENTS})255else()256 set(OpenMPTarget_FINDLIST NVPTX)257endif()258 259unset(_OpenMPTarget_MIN_VERSION)260 261# Attempt to find each requested device.262foreach(LANG IN ITEMS C CXX)263 foreach(DEVICE IN LISTS OpenMPTarget_FINDLIST)264 if(CMAKE_${LANG}_COMPILER_LOADED)265 set(OpenMPTarget_${DEVICE}_VERSION "${OpenMP_${LANG}_VERSION}")266 set(OpenMPTarget_${DEVICE}_VERSION_MAJOR "${OpenMP_${LANG}_VERSION}_MAJOR")267 set(OpenMPTarget_${DEVICE}_VERSION_MINOR "${OpenMP_${LANG}_VERSION}_MINOR")268 set(OpenMPTarget_${DEVICE}_FIND_QUIETLY ${OpenMPTarget_FIND_QUIETLY})269 set(OpenMPTarget_${DEVICE}_FIND_REQUIRED ${OpenMPTarget_FIND_REQUIRED})270 set(OpenMPTarget_${DEVICE}_FIND_VERSION ${OpenMPTarget_FIND_VERSION})271 set(OpenMPTarget_${DEVICE}_FIND_VERSION_EXACT ${OpenMPTarget_FIND_VERSION_EXACT})272 273 # OpenMP target offloading is only supported in OpenMP 4.0 an newer.274 if(OpenMPTarget_${DEVICE}_VERSION AND ("${OpenMPTarget_${DEVICE}_VERSION}" VERSION_LESS "4.0"))275 message(SEND_ERROR "FindOpenMPTarget requires at least OpenMP 4.0")276 endif()277 278 set(FPHSA_NAME_MISMATCHED TRUE)279 find_package_handle_standard_args(OpenMPTarget_${DEVICE}280 REQUIRED_VARS OpenMPTarget_${DEVICE}_FLAGS281 VERSION_VAR OpenMPTarget_${DEVICE}_VERSION)282 283 if(OpenMPTarget_${DEVICE}_FOUND)284 if(DEFINED OpenMPTarget_${DEVICE}_VERSION)285 if(NOT _OpenMPTarget_MIN_VERSION OR _OpenMPTarget_MIN_VERSION VERSION_GREATER OpenMPTarget_${LANG}_VERSION)286 set(_OpenMPTarget_MIN_VERSION OpenMPTarget_${DEVICE}_VERSION)287 endif()288 endif()289 # Create a new target.290 if(NOT TARGET OpenMPTarget::OpenMPTarget_${DEVICE})291 add_library(OpenMPTarget::OpenMPTarget_${DEVICE} INTERFACE IMPORTED)292 endif()293 # Get compiler flags for offloading to the device and architecture and294 # set the target features. Include the normal OpenMP flags as well.295 set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY296 INTERFACE_COMPILE_OPTIONS 297 "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_FLAGS}>"298 "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_ARCH}>"299 "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMP_${LANG}_FLAGS}>")300 set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY301 INTERFACE_INCLUDE_DIRECTORIES "$<BUILD_INTERFACE:${OpenMP_${LANG}_INCLUDE_DIRS}>")302 set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY303 INTERFACE_LINK_LIBRARIES 304 "${OpenMPTarget_${DEVICE}_LIBRARIES}"305 "${OpenMP_${LANG}_LIBRARIES}")306 # The offloading flags must also be passed during the linking phase so307 # the compiler can pass the binary to the correct toolchain.308 set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY309 INTERFACE_LINK_OPTIONS 310 "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_FLAGS}>"311 "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_ARCH}>"312 "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMP_${LANG}_FLAGS}>")313 # Combine all the flags if not using the target for convenience.314 set(OpenMPTarget_${DEVICE}_FLAGS ${OpenMP_${LANG}_FLAGS}315 ${OpenMPTarget_${DEVICE}_FLAGS} 316 ${OpenMPTarget_${DEVICE}_ARCH}317 CACHE STRING "${DEVICE} target compile flags for OpenMP target offloading" FORCE)318 endif()319 endif()320 endforeach()321endforeach()322 323unset(_OpenMPTarget_REQ_VARS)324foreach(DEVICE IN LISTS OpenMPTarget_FINDLIST)325 list(APPEND _OpenMPTarget_REQ_VARS "OpenMPTarget_${DEVICE}_FOUND")326endforeach()327 328find_package_handle_standard_args(OpenMPTarget329 REQUIRED_VARS ${_OpenMPTarget_REQ_VARS}330 VERSION_VAR ${_OpenMPTarget_MIN_VERSION}331 HANDLE_COMPONENTS)332 333if(NOT (CMAKE_C_COMPILER_LOADED OR CMAKE_CXX_COMPILER_LOADED) OR CMAKE_Fortran_COMPILER_LOADED)334 message(SEND_ERROR "FindOpenMPTarget requires the C or CXX languages to be enabled")335endif()336 337unset(OpenMPTarget_C_CXX_TEST_SOURCE)338cmake_policy(POP)339