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-exception
6#
7##===----------------------------------------------------------------------===##
8#
9# Find OpenMP Target offloading Support for various compilers.
10#
11##===----------------------------------------------------------------------===##
12
13#[========================================================================[.rst:
14FindOpenMPTarget
15----------------
16
17Finds OpenMP Target Offloading Support.
18
19This module can be used to detect OpenMP target offloading support in a
20compiler. If the compiler support OpenMP Offloading to a specified target, the
21flags required to compile offloading code to that target are output for each
22target.
23
24This module will automatically include OpenMP support if it was not loaded
25already. It does not need to be included separately to get full OpenMP support.
26
27Variables
28^^^^^^^^^
29
30The module exposes the components ``NVPTX`` and ``AMDGCN``.  Each of these
31controls the various offloading targets to search OpenMP target offloasing
32support 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 requested
38  targets have been found.
39
40This module will set the following variables per language in your
41project, where ``<device>`` is one of NVPTX or AMDGCN
42
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 following
49variables 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 manually
60specified by setting a variable to the desired architecture. Variables can also
61be 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 NVPTX
65  or `gfx908` for AMDGCN.
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 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
80
81cmake_policy(PUSH)
82cmake_policy(VERSION 3.13.4)
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 "AMDGCN")
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  # AMD requires the architecture, default to gfx908 if not provided.
144  if((NOT OpenMPTarget_${DEVICE}_ARCH) AND ("${DEVICE}" STREQUAL "AMDGCN"))
145    set(OpenMPTarget_${DEVICE}_ARCH "gfx908")
146  endif()
147  if(OpenMPTarget_${DEVICE}_ARCH)
148    # Only Clang supports selecting the architecture for now.
149    set(OMPTarget_ARCH_Clang "-Xopenmp-target=${DEVICE_FLAG} -march=${OpenMPTarget_${DEVICE}_ARCH}")
150
151    if(DEFINED OMPTarget_ARCH_${CMAKE_${LANG}_COMPILER_ID})
152      set(OpenMPTarget_DEVICE_ARCH_CANDIDATES "${OMPTarget_ARCH_${CMAKE_${LANG}_COMPILER_ID}}")
153    endif()
154    set(OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES "${OpenMPTarget_DEVICE_ARCH_CANDIDATES}" PARENT_SCOPE)
155  else()
156    set(OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES "" PARENT_SCOPE)
157  endif()
158endfunction()
159
160set(OpenMPTarget_C_CXX_TEST_SOURCE
161"#include <omp.h>
162int main(void) {
163  int isHost;
164#pragma omp target map(from: isHost)
165  { isHost = omp_is_initial_device(); }
166  return isHost;
167}")
168
169function(_OPENMP_TARGET_WRITE_SOURCE_FILE LANG SRC_FILE_CONTENT_VAR SRC_FILE_NAME SRC_FILE_FULLPATH)
170  set(WORK_DIR ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/FindOpenMPTarget)
171  if("${LANG}" STREQUAL "C")
172    set(SRC_FILE "${WORK_DIR}/${SRC_FILE_NAME}.c")
173    file(WRITE "${SRC_FILE}" "${OpenMPTarget_C_CXX_${SRC_FILE_CONTENT_VAR}}")
174  elseif("${LANG}" STREQUAL "CXX")
175    set(SRC_FILE "${WORK_DIR}/${SRC_FILE_NAME}.cpp")
176    file(WRITE "${SRC_FILE}" "${OpenMPTarget_C_CXX_${SRC_FILE_CONTENT_VAR}}")
177  endif()
178  set(${SRC_FILE_FULLPATH} "${SRC_FILE}" PARENT_SCOPE)
179endfunction()
180
181# Get the candidate flags and try to compile a test application. If it compiles
182# and all the flags are found, we assume the compiler supports offloading.
183function(_OPENMP_TARGET_DEVICE_GET_FLAGS LANG DEVICE OPENMP_FLAG_VAR OPENMP_LIB_VAR OPENMP_DEVICE_VAR OPENMP_ARCH_VAR)
184  _OPENMP_TARGET_DEVICE_CANDIDATES(${LANG} ${DEVICE})
185  _OPENMP_TARGET_DEVICE_FLAG_CANDIDATES(${LANG} "${OpenMPTarget_${LANG}_DEVICE_CANDIDATES}")
186  _OPENMP_TARGET_DEVICE_ARCH_CANDIDATES(${LANG} ${DEVICE} "${OpenMPTarget_${LANG}_DEVICE_CANDIDATES}")
187  _OPENMP_TARGET_WRITE_SOURCE_FILE("${LANG}" "TEST_SOURCE" OpenMPTargetTryFlag _OPENMP_TEST_SRC)
188
189  # Try to compile a test application with the found flags.
190  try_compile(OpenMPTarget_COMPILE_RESULT ${CMAKE_BINARY_DIR} ${_OPENMP_TEST_SRC}
191    CMAKE_FLAGS "-DCOMPILE_DEFINITIONS:STRING=${OpenMP_${LANG}_FLAGS} ${OpenMPTarget_${LANG}_FLAG_CANDIDATES} ${OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES}"
192    "-DINCLUDE_DIRECTORIES:STRING=${OpenMP_${LANG}_INCLUDE_DIR}"
193    LINK_LIBRARIES ${CMAKE_${LANG}_VERBOSE_FLAG}
194    OUTPUT_VARIABLE OpenMP_TRY_COMPILE_OUTPUT
195  )
196
197  file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log
198    "Detecting OpenMP ${CMAKE_${LANG}_COMPILER_ID} ${DEVICE} target support with the following Flags:
199    ${OpenMP_${LANG}_FLAGS} ${OpenMPTarget_${LANG}_FLAG_CANDIDATES} ${OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES}
200    With the following output:\n ${OpenMP_TRY_COMPILE_OUTPUT}\n")
201
202  # If compilation was successful and the device was found set the return variables.
203  if (OpenMPTarget_COMPILE_RESULT AND DEFINED OpenMPTarget_${LANG}_DEVICE_CANDIDATES)
204    file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log
205      "Compilation successful, adding flags for ${DEVICE}.\n\n")
206
207    # Clang has a seperate library for target offloading.
208    if(CMAKE_${LANG}_COMPILER_ID STREQUAL "Clang")
209      find_library(OpenMPTarget_libomptarget_LIBRARY
210        NAMES omptarget
211        HINTS ${CMAKE_${LANG}_IMPLICIT_LINK_DIRECTORIES}
212      )
213      mark_as_advanced(OpenMPTarget_libomptarget_LIBRARY)
214      set("${OPENMP_LIB_VAR}" "${OpenMPTarget_libomptarget_LIBRARY}" PARENT_SCOPE)
215    else()
216      unset("${OPENMP_LIB_VAR}" PARENT_SCOPE)
217    endif()
218    set("${OPENMP_DEVICE_VAR}" "${OpenMPTarget_${LANG}_DEVICE_CANDIDATES}" PARENT_SCOPE)
219    set("${OPENMP_FLAG_VAR}" "${OpenMPTarget_${LANG}_FLAG_CANDIDATES}" PARENT_SCOPE)
220    set("${OPENMP_ARCH_VAR}" "${OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES}" PARENT_SCOPE)
221  else()
222    unset("${OPENMP_DEVICE_VAR}" PARENT_SCOPE)
223    unset("${OPENMP_FLAG_VAR}" PARENT_SCOPE)
224    unset("${OPENMP_ARCH_VAR}" PARENT_SCOPE)
225  endif()
226endfunction()
227
228# Load the compiler support for each device.
229foreach(LANG IN ITEMS C CXX)
230  # Cache the version in case CMake doesn't load the OpenMP package this time
231  set(OpenMP_${LANG}_VERSION ${OpenMP_${LANG}_VERSION}
232    CACHE STRING "OpenMP Version" FORCE)
233  mark_as_advanced(OpenMP_${LANG}_VERSION)
234  foreach(DEVICE IN ITEMS NVPTX AMDGCN)
235    if(CMAKE_${LANG}_COMPILER_LOADED)
236      if(NOT DEFINED OpenMPTarget_${LANG}_FLAGS OR NOT DEFINED OpenMPTarget_${LANG}_DEVICE)
237        _OPENMP_TARGET_DEVICE_GET_FLAGS(${LANG} ${DEVICE}
238          OpenMPTarget_${DEVICE}_FLAGS_WORK
239          OpenMPTarget_${DEVICE}_LIBS_WORK
240          OpenMPTarget_${DEVICE}_DEVICE_WORK
241          OpenMPTarget_${DEVICE}_ARCHS_WORK)
242
243        separate_arguments(_OpenMPTarget_${DEVICE}_FLAGS NATIVE_COMMAND "${OpenMPTarget_${DEVICE}_FLAGS_WORK}")
244        separate_arguments(_OpenMPTarget_${DEVICE}_ARCHS NATIVE_COMMAND "${OpenMPTarget_${DEVICE}_ARCHS_WORK}")
245        set(OpenMPTarget_${DEVICE}_FLAGS ${_OpenMPTarget_${DEVICE}_FLAGS}
246            CACHE STRING "${DEVICE} target compile flags for OpenMP target offloading" FORCE)
247        set(OpenMPTarget_${DEVICE}_ARCH ${_OpenMPTarget_${DEVICE}_ARCHS}
248            CACHE STRING "${DEVICE} target architecture flags for OpenMP target offloading" FORCE)
249        set(OpenMPTarget_${DEVICE}_LIBS ${OpenMPTarget_${DEVICE}_LIBS_WORK}
250            CACHE STRING "${DEVICE} target libraries for OpenMP target offloading" FORCE)
251        mark_as_advanced(OpenMPTarget_${DEVICE}_FLAGS OpenMPTarget_${DEVICE}_ARCH OpenMPTarget_${DEVICE}_LIBS)
252      endif()
253    endif()
254  endforeach()
255endforeach()
256
257if(OpenMPTarget_FIND_COMPONENTS)
258  set(OpenMPTarget_FINDLIST ${OpenMPTarget_FIND_COMPONENTS})
259else()
260  set(OpenMPTarget_FINDLIST NVPTX)
261endif()
262
263unset(_OpenMPTarget_MIN_VERSION)
264
265# Attempt to find each requested device.
266foreach(LANG IN ITEMS C CXX)
267  foreach(DEVICE IN LISTS OpenMPTarget_FINDLIST)
268    if(CMAKE_${LANG}_COMPILER_LOADED)
269      set(OpenMPTarget_${DEVICE}_VERSION "${OpenMP_${LANG}_VERSION}")
270      set(OpenMPTarget_${DEVICE}_VERSION_MAJOR "${OpenMP_${LANG}_VERSION}_MAJOR")
271      set(OpenMPTarget_${DEVICE}_VERSION_MINOR "${OpenMP_${LANG}_VERSION}_MINOR")
272      set(OpenMPTarget_${DEVICE}_FIND_QUIETLY ${OpenMPTarget_FIND_QUIETLY})
273      set(OpenMPTarget_${DEVICE}_FIND_REQUIRED ${OpenMPTarget_FIND_REQUIRED})
274      set(OpenMPTarget_${DEVICE}_FIND_VERSION ${OpenMPTarget_FIND_VERSION})
275      set(OpenMPTarget_${DEVICE}_FIND_VERSION_EXACT ${OpenMPTarget_FIND_VERSION_EXACT})
276
277      # OpenMP target offloading is only supported in OpenMP 4.0 an newer.
278      if(OpenMPTarget_${DEVICE}_VERSION AND ("${OpenMPTarget_${DEVICE}_VERSION}" VERSION_LESS "4.0"))
279        message(SEND_ERROR "FindOpenMPTarget requires at least OpenMP 4.0")
280      endif()
281
282      set(FPHSA_NAME_MISMATCHED TRUE)
283      find_package_handle_standard_args(OpenMPTarget_${DEVICE}
284        REQUIRED_VARS OpenMPTarget_${DEVICE}_FLAGS
285        VERSION_VAR OpenMPTarget_${DEVICE}_VERSION)
286
287      if(OpenMPTarget_${DEVICE}_FOUND)
288        if(DEFINED OpenMPTarget_${DEVICE}_VERSION)
289          if(NOT _OpenMPTarget_MIN_VERSION OR _OpenMPTarget_MIN_VERSION VERSION_GREATER OpenMPTarget_${LANG}_VERSION)
290            set(_OpenMPTarget_MIN_VERSION OpenMPTarget_${DEVICE}_VERSION)
291          endif()
292        endif()
293        # Create a new target.
294        if(NOT TARGET OpenMPTarget::OpenMPTarget_${DEVICE})
295          add_library(OpenMPTarget::OpenMPTarget_${DEVICE} INTERFACE IMPORTED)
296        endif()
297        # Get compiler flags for offloading to the device and architecture and
298        # set the target features. Include the normal OpenMP flags as well.
299        set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
300          INTERFACE_COMPILE_OPTIONS
301          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_FLAGS}>"
302          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_ARCH}>"
303          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMP_${LANG}_FLAGS}>")
304        set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
305          INTERFACE_INCLUDE_DIRECTORIES "$<BUILD_INTERFACE:${OpenMP_${LANG}_INCLUDE_DIRS}>")
306        set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
307          INTERFACE_LINK_LIBRARIES
308          "${OpenMPTarget_${DEVICE}_LIBS}"
309          "${OpenMP_${LANG}_LIBRARIES}")
310        # The offloading flags must also be passed during the linking phase so
311        # the compiler can pass the binary to the correct toolchain.
312        set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
313          INTERFACE_LINK_OPTIONS
314          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_FLAGS}>"
315          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_ARCH}>"
316          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMP_${LANG}_FLAGS}>")
317        # Combine all the flags if not using the target for convenience.
318        set(OpenMPTarget_${DEVICE}_FLAGS ${OpenMP_${LANG}_FLAGS}
319          ${OpenMPTarget_${DEVICE}_FLAGS}
320          ${OpenMPTarget_${DEVICE}_ARCH}
321          CACHE STRING "${DEVICE} target compile flags for OpenMP target offloading" FORCE)
322      endif()
323    endif()
324  endforeach()
325endforeach()
326
327unset(_OpenMPTarget_REQ_VARS)
328foreach(DEVICE IN LISTS OpenMPTarget_FINDLIST)
329  list(APPEND _OpenMPTarget_REQ_VARS "OpenMPTarget_${DEVICE}_FOUND")
330endforeach()
331
332find_package_handle_standard_args(OpenMPTarget
333    REQUIRED_VARS ${_OpenMPTarget_REQ_VARS}
334    VERSION_VAR ${_OpenMPTarget_MIN_VERSION}
335    HANDLE_COMPONENTS)
336
337if(NOT (CMAKE_C_COMPILER_LOADED OR CMAKE_CXX_COMPILER_LOADED) OR CMAKE_Fortran_COMPILER_LOADED)
338  message(SEND_ERROR "FindOpenMPTarget requires the C or CXX languages to be enabled")
339endif()
340
341unset(OpenMPTarget_C_CXX_TEST_SOURCE)
342cmake_policy(POP)
343