1 ##===----------------------------------------------------------------------===##
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-exception
7 ##===----------------------------------------------------------------------===##
9 # Find OpenMP Target offloading Support for various compilers.
11 ##===----------------------------------------------------------------------===##
13 #[========================================================================[.rst:
17 Finds OpenMP Target Offloading Support.
19 This module can be used to detect OpenMP target offloading support in a
20 compiler. If the compiler support OpenMP Offloading to a specified target, the
21 flags required to compile offloading code to that target are output for each
24 This module will automatically include OpenMP support if it was not loaded
25 already. It does not need to be included separately to get full OpenMP support.
30 The module exposes the components ``NVPTX`` and ``AMDGPU``. Each of these
31 controls the various offloading targets to search OpenMP target offloasing
34 Depending on the enabled components the following variables will be set:
36 ``OpenMPTarget_FOUND``
37 Variable indicating that OpenMP target offloading flags for all requested
38 targets have been found.
40 This module will set the following variables per language in your
41 project, where ``<device>`` is one of NVPTX or AMDGPU
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.
48 For linking with OpenMP code written in ``<device>``, the following
49 variables are provided:
51 ``OpenMPTarget_<device>_LIBRARIES``
52 A list of libraries needed to link with OpenMP code written in ``<lang>``.
54 Additionally, the module provides :prop_tgt:`IMPORTED` targets:
56 ``OpenMPTarget::OpenMPTarget_<device>``
57 Target for using OpenMP offloading to ``<device>``.
59 If the specific architecture of the target is needed, it can be manually
60 specified by setting a variable to the desired architecture. Variables can also
61 be used to override the standard flag searching for a given compiler.
63 ``OpenMPTarget_<device>_ARCH``
64 Sets the architecture of ``<device>`` to compile for. Such as `sm_70` for NVPTX
65 or `gfx908` for AMDGPU.
67 ``OpenMPTarget_<device>_DEVICE``
68 Sets the name of the device to offload to.
70 ``OpenMPTarget_<device>_FLAGS``
71 Sets the compiler flags for offloading to ``<device>``.
73 #]========================================================================]
75 # TODO: Support Fortran
76 # TODO: Support multiple offloading targets by setting the "OpenMPTarget" target
77 # to include flags for all components loaded
78 # TODO: Configure target architecture without a variable (component NVPTX_SM_70)
79 # TODO: Test more compilers
82 cmake_policy(VERSION 3.20.0)
84 find_package(OpenMP ${OpenMPTarget_FIND_VERSION} REQUIRED)
86 # Find the offloading flags for each compiler.
87 function(_OPENMP_TARGET_DEVICE_FLAG_CANDIDATES LANG DEVICE)
88 if(NOT OpenMPTarget_${LANG}_FLAGS)
89 unset(OpenMPTarget_FLAG_CANDIDATES)
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}")
97 if(DEFINED OMPTarget_FLAGS_${CMAKE_${LANG}_COMPILER_ID})
98 set(OpenMPTarget_FLAG_CANDIDATES "${OMPTarget_FLAGS_${CMAKE_${LANG}_COMPILER_ID}}")
101 set(OpenMPTarget_${LANG}_FLAG_CANDIDATES "${OpenMPTarget_FLAG_CANDIDATES}" PARENT_SCOPE)
103 set(OpenMPTarget_${LANG}_FLAG_CANDIDATES "${OpenMPTarget_${LANG}_FLAGS}" PARENT_SCOPE)
107 # Get the coded name of the device for each compiler.
108 function(_OPENMP_TARGET_DEVICE_CANDIDATES LANG DEVICE)
109 if (NOT OpenMPTarget_${DEVICE}_DEVICE)
110 unset(OpenMPTarget_DEVICE_CANDIDATES)
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")
117 set(OMPTarget_DEVICE_Clang "nvptx64-nvidia-cuda")
119 set(OMPTarget_DEVICE_GNU "nvptx-none")
120 set(OMPTarget_DEVICE_XL "")
121 set(OMPTarget_DEVICE_PGI "gpu")
122 set(OMPTarget_DEVICE_NVHPC "gpu")
124 if(DEFINED OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID})
125 set(OpenMPTarget_DEVICE_CANDIDATES "${OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID}}")
127 elseif("${DEVICE}" STREQUAL "AMDGPU")
128 set(OMPTarget_DEVICE_Clang "amdgcn-amd-amdhsa")
129 set(OMPTarget_DEVICE_GNU "hsa")
131 if(DEFINED OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID})
132 set(OpenMPTarget_DEVICE_CANDIDATES "${OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID}}")
135 set(OpenMPTarget_${LANG}_DEVICE_CANDIDATES "${OpenMPTarget_DEVICE_CANDIDATES}" PARENT_SCOPE)
137 set(OpenMPTarget_${LANG}_DEVICE_CANDIDATES "${OpenMPTarget_${LANG}_DEVICE}" PARENT_SCOPE)
141 # Get flags for setting the device's architecture for each compiler.
142 function(_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}")
147 if(DEFINED OMPTarget_ARCH_${CMAKE_${LANG}_COMPILER_ID})
148 set(OpenMPTarget_DEVICE_ARCH_CANDIDATES "${OMPTarget_ARCH_${CMAKE_${LANG}_COMPILER_ID}}")
150 set(OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES "${OpenMPTarget_DEVICE_ARCH_CANDIDATES}" PARENT_SCOPE)
152 set(OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES "" PARENT_SCOPE)
156 set(OpenMPTarget_C_CXX_TEST_SOURCE
160 #pragma omp target map(from: isHost)
161 { isHost = omp_is_initial_device(); }
165 function(_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}}")
174 set(${SRC_FILE_FULLPATH} "${SRC_FILE}" PARENT_SCOPE)
177 # Get the candidate flags and try to compile a test application. If it compiles
178 # and all the flags are found, we assume the compiler supports offloading.
179 function(_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)
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_OUTPUT
193 file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log
194 "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")
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.log
201 "Compilation successful, adding flags for ${DEVICE}.\n\n")
203 # Clang has a separate library for target offloading.
204 if(CMAKE_${LANG}_COMPILER_ID STREQUAL "Clang")
205 find_library(OpenMPTarget_libomptarget_LIBRARY
207 HINTS ${CMAKE_${LANG}_IMPLICIT_LINK_DIRECTORIES}
209 mark_as_advanced(OpenMPTarget_libomptarget_LIBRARY)
210 set("${OPENMP_LIB_VAR}" "${OpenMPTarget_libomptarget_LIBRARY}" PARENT_SCOPE)
212 unset("${OPENMP_LIB_VAR}" PARENT_SCOPE)
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)
218 unset("${OPENMP_DEVICE_VAR}" PARENT_SCOPE)
219 unset("${OPENMP_FLAG_VAR}" PARENT_SCOPE)
220 unset("${OPENMP_ARCH_VAR}" PARENT_SCOPE)
224 # Load the compiler support for each device.
225 foreach(LANG IN ITEMS C CXX)
226 # Cache the version in case CMake doesn't load the OpenMP package this time
227 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_WORK
235 OpenMPTarget_${DEVICE}_LIBS_WORK
236 OpenMPTarget_${DEVICE}_DEVICE_WORK
237 OpenMPTarget_${DEVICE}_ARCHS_WORK)
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)
253 if(OpenMPTarget_FIND_COMPONENTS)
254 set(OpenMPTarget_FINDLIST ${OpenMPTarget_FIND_COMPONENTS})
256 set(OpenMPTarget_FINDLIST NVPTX)
259 unset(_OpenMPTarget_MIN_VERSION)
261 # Attempt to find each requested device.
262 foreach(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})
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")
278 set(FPHSA_NAME_MISMATCHED TRUE)
279 find_package_handle_standard_args(OpenMPTarget_${DEVICE}
280 REQUIRED_VARS OpenMPTarget_${DEVICE}_FLAGS
281 VERSION_VAR OpenMPTarget_${DEVICE}_VERSION)
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)
289 # Create a new target.
290 if(NOT TARGET OpenMPTarget::OpenMPTarget_${DEVICE})
291 add_library(OpenMPTarget::OpenMPTarget_${DEVICE} INTERFACE IMPORTED)
293 # Get compiler flags for offloading to the device and architecture and
294 # set the target features. Include the normal OpenMP flags as well.
295 set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
296 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} PROPERTY
301 INTERFACE_INCLUDE_DIRECTORIES "$<BUILD_INTERFACE:${OpenMP_${LANG}_INCLUDE_DIRS}>")
302 set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
303 INTERFACE_LINK_LIBRARIES
304 "${OpenMPTarget_${DEVICE}_LIBRARIES}"
305 "${OpenMP_${LANG}_LIBRARIES}")
306 # The offloading flags must also be passed during the linking phase so
307 # the compiler can pass the binary to the correct toolchain.
308 set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
309 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)
323 unset(_OpenMPTarget_REQ_VARS)
324 foreach(DEVICE IN LISTS OpenMPTarget_FINDLIST)
325 list(APPEND _OpenMPTarget_REQ_VARS "OpenMPTarget_${DEVICE}_FOUND")
328 find_package_handle_standard_args(OpenMPTarget
329 REQUIRED_VARS ${_OpenMPTarget_REQ_VARS}
330 VERSION_VAR ${_OpenMPTarget_MIN_VERSION}
333 if(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")
337 unset(OpenMPTarget_C_CXX_TEST_SOURCE)