[OpenMP] Introduce an CMake find module for OpenMP Target support

[jhuber6]

Summary

This introduces a CMake find module for detecting target offloading support in
a compiler. The goal is to make it easier to incorporate target offloading into
a cmake project.

[tianshilei1992]

LGTM. Some nits.

openmp/tools/Modules/FindOpenMPTarget.cmake

2

Is it better to put these files into libomptarget?

88

Not sure if it is too strict. LLVM and OpenMP both require CMake 3.13.4 and above.

[jhuber6]

openmp/tools/Modules/FindOpenMPTarget.cmake

2

We don't have a tools directory under that, I figured it's fine in the same way we put the libomptarget documentation under openmp/docs.

88

I think the only things preventing me from using 3.13.4 is that the OpenMP module doesn't seem to return the version correctly, and it doesn't support the "NAME_MISMATCHED" which means that newer CMake will emit a warning for finding a package with a different name than the one specified.

[tianshilei1992]

openmp/tools/Modules/FindOpenMPTarget.cmake

88

The feature is good for sure, but since we require only 3.13+ to build OpenMP, I think it's better to keep it consistent; otherwise it might be confusing.

[ye-luo]

Is this intended for libomptarget only or even to take care other offload runtime libraries?
if it is the former case, It needs to be renamed as FindLibOMPTarget to be specific about libomptarget.
If it is the later case, please upstream it to CMake directly.

User code may want to defined either own OpenMPTarget target and alias it to the target created by this module if they use libomptarget or define others if other vendor OpenMP offload runtime library is used.

OpenMP is both a language extension and library, creating a cmake module is a challenge.

openmp/tools/Modules/FindOpenMPTarget.cmake

63

I don't quite follow what is OpenMP version especially the device specific one.
As vendors usually only implement a subset, users should not rely on OpenMP spec version anyway.
If they really need it, just invoke FindOpenMP themselves. RIght now it is an illusion there is a separate OpenMPTarget_VERSION

100

FLAGS, plural because you might need multiple flags.

Use NVHPC instead of PGI, the correct flag is -mp=gpu
does -mp=${DEVICE} work

GNU offload to NVIDIA is -foffload=nvptx-none
does -foffload=${DEVICE} work?

[jhuber6]

Changing the minimum version requirements.

[jhuber6]

In D104710#2833720, @ye-luo wrote:

Is this intended for libomptarget only or even to take care other offload runtime libraries?
if it is the former case, It needs to be renamed as FindLibOMPTarget to be specific about libomptarget.
If it is the later case, please upstream it to CMake directly.

I intend for this to work for other compilers, but I don't have as much experience with the other compilers so I can't guarantee its effectiveness. Upstreaming to CMake could be a goal, but I think for now it's easier to keep it here until some more people have has a chance to test it.

User code may want to defined either own OpenMPTarget target and alias it to the target created by this module if they use libomptarget or define others if other vendor OpenMP offload runtime library is used.

You should be able to manually set the flags by setting the OpenMPTarget_<Device>_FLAGS or OpenMPTarget_<Device>_DEVICE since it will forward those if they're already set.

OpenMP is both a language extension and library, creating a cmake module is a challenge.

Yeah, I primarily copied the method used for the existing FindOpenMP.cmake module, fortunately I can reuse a lot of features.

openmp/tools/Modules/FindOpenMPTarget.cmake

63

Since FindOpenMP returns a version, I think it's reasonable to return the same version for this module because it's intended to be a superset. It should include all the varialbes FindOpenMP offers, plus the ones needed for offloading. (I may need to check if I set the include dirs and libs correctly in addition).

100

PGI what the COMPILDER_ID returns when I ran with the NVHPC compiler last I checked. I think it only supports nvptx so there's no option to configure its target.

GNU allows this flag, technically linking math and atomic on the device are optional, but they are needed in most applications so it's much easier to do it by default.

I'll make it plural.

[ye-luo]

In D104710#2833784, @jhuber6 wrote:

In D104710#2833720, @ye-luo wrote:

Is this intended for libomptarget only or even to take care other offload runtime libraries?
if it is the former case, It needs to be renamed as FindLibOMPTarget to be specific about libomptarget.
If it is the later case, please upstream it to CMake directly.

I intend for this to work for other compilers, but I don't have as much experience with the other compilers so I can't guarantee its effectiveness. Upstreaming to CMake could be a goal, but I think for now it's easier to keep it here until some more people have has a chance to test it.

other compiler with libomp+libomptarget or other compilers with their own OpenMP runtime libraries?

openmp/tools/Modules/FindOpenMPTarget.cmake

63

I mean it is better not to add another version variable if the OpenMP_VERSION is sufficient and has been included already.

100

3.20 changed PGI to NVHPC

[jhuber6]

Other compilers, Clang is the only one with a separate library AFAIK. The others just put them inside of the existing ones like libgomp so I don't search for them.

openmp/tools/Modules/FindOpenMPTarget.cmake

63

I could get rid of it, or just unset them manually. I mostly just wanted the versions to make the find_package_handle_standard_args work more in line with OpenMP.

100

Will probably need to provide both since this should work for older versions as well.

[jhuber6]

Addressing comments.

[jhuber6]

Fixing small error. I've checked gcc, clang, xlc, and nvc offloading to NVPTX
and only clang offloading to AMDGCN and it seems to work for those.

openmp/tools/Modules/CMakeLists.txt

+##===----------------------------------------------------------------------===##
+#
+# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#
+##===----------------------------------------------------------------------===##
+#
+# Install the FindOpenMPTarget module
+#
+##===----------------------------------------------------------------------===##
+
+
+install(FILES "FindOpenMPTarget.cmake"
+              DESTINATION "${OPENMP_INSTALL_LIBDIR}/cmake/openmp")

openmp/tools/Modules/FindOpenMPTarget.cmake

+##===----------------------------------------------------------------------===##
+#
+# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#
+##===----------------------------------------------------------------------===##
+#
+# Find OpenMP Target offloading Support for various compilers.
+#
+##===----------------------------------------------------------------------===##
+
+#[========================================================================[.rst:
+FindOpenMPTarget
+----------------
+
+Finds OpenMP Target Offloading Support.
+
+This module can be used to detect OpenMP target offloading support in a
+compiler. If the compiler support OpenMP Offloading to a specified target, the
+flags required to compile offloading code to that target are output for each
+target.
+
+This module will automatically include OpenMP support if it was not loaded
+already. It does not need to be included separately to get full OpenMP support.
+
+Variables
+^^^^^^^^^
+
+The module exposes the components ``NVPTX`` and ``AMDGCN``.  Each of these
+controls the various offloading targets to search OpenMP target offloasing
+support for.
+
+Depending on the enabled components the following variables will be set:
+
+``OpenMPTarget_FOUND``
+  Variable indicating that OpenMP target offloading flags for all requested
+  targets have been found.
+
+This module will set the following variables per language in your
+project, where ``<device>`` is one of NVPTX or AMDGCN
+
+``OpenMPTarget_<device>_FOUND``
+  Variable indicating if OpenMP support for the ``<device>`` was detected.
+``OpenMPTarget_<device>_FLAGS``
+  OpenMP compiler flags for offloading to ``<device>``, separated by spaces.
+
+For linking with OpenMP code written in ``<device>``, the following
+variables are provided:
+
+``OpenMPTarget_<device>_LIBRARIES``
+  A list of libraries needed to link with OpenMP code written in ``<lang>``.
+
+Additionally, the module provides :prop_tgt:`IMPORTED` targets:
+
+``OpenMPTarget::OpenMPTarget_<device>``
+  Target for using OpenMP offloading to ``<device>``.
+
+If the specific architecture of the target is needed, it can be manually
+specified by setting a variable to the desired architecture. Variables can also
+be used to override the standard flag searching for a given compiler.
+
+``OpenMPTarget_<device>_ARCH``
+  Sets the architecture of ``<device>`` to compile for. Such as `sm_70` for NVPTX
+  or `gfx908` for AMDGCN. 
+
+``OpenMPTarget_<device>_DEVICE``
+  Sets the name of the device to offload to.
+
+``OpenMPTarget_<device>_FLAGS``
+  Sets the compiler flags for offloading to ``<device>``.
+
+#]========================================================================]
+
+# TODO: Support Fortran
+# TODO: Support multiple offloading targets by setting the "OpenMPTarget" target
+#       to include flags for all components loaded
+# TODO: Configure target architecture without a variable (component NVPTX_SM_70)
+# TODO: Test more compilers
+
+cmake_policy(PUSH)
+cmake_policy(VERSION 3.13.4)
+
+find_package(OpenMP ${OpenMPTarget_FIND_VERSION} REQUIRED)
+
+# Find the offloading flags for each compiler.
+function(_OPENMP_TARGET_DEVICE_FLAG_CANDIDATES LANG DEVICE)
+  if(NOT OpenMPTarget_${LANG}_FLAGS)
+    unset(OpenMPTarget_FLAG_CANDIDATES)
+
+    set(OMPTarget_FLAGS_Clang "-fopenmp-targets=${DEVICE}")
+    set(OMPTarget_FLAGS_GNU "-foffload=${DEVICE}=\"-lm -latomic\"")
+    set(OMPTarget_FLAGS_XL "-qoffload")
+    set(OMPTarget_FLAGS_PGI "-mp=${DEVICE}")
+    set(OMPTarget_FLAGS_NVHPC "-mp=${DEVICE}")
+
+    if(DEFINED OMPTarget_FLAGS_${CMAKE_${LANG}_COMPILER_ID})
+      set(OpenMPTarget_FLAG_CANDIDATES "${OMPTarget_FLAGS_${CMAKE_${LANG}_COMPILER_ID}}")
+    endif()
+
+    set(OpenMPTarget_${LANG}_FLAG_CANDIDATES "${OpenMPTarget_FLAG_CANDIDATES}" PARENT_SCOPE)
+  else()
+    set(OpenMPTarget_${LANG}_FLAG_CANDIDATES "${OpenMPTarget_${LANG}_FLAGS}" PARENT_SCOPE)
+  endif()
+endfunction()
+
+# Get the coded name of the device for each compiler.
+function(_OPENMP_TARGET_DEVICE_CANDIDATES LANG DEVICE)
+  if (NOT OpenMPTarget_${DEVICE}_DEVICE)
+    unset(OpenMPTarget_DEVICE_CANDIDATES)
+
+    # Check each supported device.
+    if("${DEVICE}" STREQUAL "NVPTX")
+      if ("${CMAKE_SIZEOF_VOID_P}" STREQUAL "4")
+        set(OMPTarget_DEVICE_Clang "nvptx32-nvidia-cuda")
+      else()
+        set(OMPTarget_DEVICE_Clang "nvptx64-nvidia-cuda")
+      endif()
+      set(OMPTarget_DEVICE_GNU "nvptx-none")
+      set(OMPTarget_DEVICE_XL "")
+      set(OMPTarget_DEVICE_PGI "gpu")
+      set(OMPTarget_DEVICE_NVHPC "gpu")
+
+      if(DEFINED OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID})
+        set(OpenMPTarget_DEVICE_CANDIDATES "${OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID}}")
+      endif()
+    elseif("${DEVICE}" STREQUAL "AMDGCN")
+      set(OMPTarget_DEVICE_Clang "amdgcn-amd-amdhsa")
+      set(OMPTarget_DEVICE_GNU "hsa")
+
+      if(DEFINED OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID})
+        set(OpenMPTarget_DEVICE_CANDIDATES "${OMPTarget_DEVICE_${CMAKE_${LANG}_COMPILER_ID}}")
+      endif()
+    endif()
+    set(OpenMPTarget_${LANG}_DEVICE_CANDIDATES "${OpenMPTarget_DEVICE_CANDIDATES}" PARENT_SCOPE)
+  else()
+    set(OpenMPTarget_${LANG}_DEVICE_CANDIDATES "${OpenMPTarget_${LANG}_DEVICE}" PARENT_SCOPE)
+  endif()
+endfunction()
+
+# Get flags for setting the device's architecture for each compiler.
+function(_OPENMP_TARGET_DEVICE_ARCH_CANDIDATES LANG DEVICE DEVICE_FLAG)
+  # AMD requires the architecture, default to gfx908 if not provided.
+  if((NOT OpenMPTarget_${DEVICE}_ARCH) AND ("${DEVICE}" STREQUAL "AMDGCN"))
+    set(OpenMPTarget_${DEVICE}_ARCH "gfx908")
+  endif()
+  if(OpenMPTarget_${DEVICE}_ARCH)
+    # Only Clang supports selecting the architecture for now.
+    set(OMPTarget_ARCH_Clang "-Xopenmp-target=${DEVICE_FLAG} -march=${OpenMPTarget_${DEVICE}_ARCH}")
+
+    if(DEFINED OMPTarget_ARCH_${CMAKE_${LANG}_COMPILER_ID})
+      set(OpenMPTarget_DEVICE_ARCH_CANDIDATES "${OMPTarget_ARCH_${CMAKE_${LANG}_COMPILER_ID}}")
+    endif()
+    set(OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES "${OpenMPTarget_DEVICE_ARCH_CANDIDATES}" PARENT_SCOPE)
+  else()
+    set(OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES "" PARENT_SCOPE)
+  endif()
+endfunction()
+
+set(OpenMPTarget_C_CXX_TEST_SOURCE
+"#include <omp.h>
+int main(void) {
+  int isHost;
+#pragma omp target map(from: isHost)
+  { isHost = omp_is_initial_device(); }
+  return isHost;
+}")
+
+function(_OPENMP_TARGET_WRITE_SOURCE_FILE LANG SRC_FILE_CONTENT_VAR SRC_FILE_NAME SRC_FILE_FULLPATH)
+  set(WORK_DIR ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/FindOpenMPTarget)
+  if("${LANG}" STREQUAL "C")
+    set(SRC_FILE "${WORK_DIR}/${SRC_FILE_NAME}.c")
+    file(WRITE "${SRC_FILE}" "${OpenMPTarget_C_CXX_${SRC_FILE_CONTENT_VAR}}")
+  elseif("${LANG}" STREQUAL "CXX")
+    set(SRC_FILE "${WORK_DIR}/${SRC_FILE_NAME}.cpp")
+    file(WRITE "${SRC_FILE}" "${OpenMPTarget_C_CXX_${SRC_FILE_CONTENT_VAR}}")
+  endif()
+  set(${SRC_FILE_FULLPATH} "${SRC_FILE}" PARENT_SCOPE)
+endfunction()
+
+# Get the candidate flags and try to compile a test application. If it compiles
+# and all the flags are found, we assume the compiler supports offloading.
+function(_OPENMP_TARGET_DEVICE_GET_FLAGS LANG DEVICE OPENMP_FLAG_VAR OPENMP_LIB_VAR OPENMP_DEVICE_VAR OPENMP_ARCH_VAR)
+  _OPENMP_TARGET_DEVICE_CANDIDATES(${LANG} ${DEVICE})
+  _OPENMP_TARGET_DEVICE_FLAG_CANDIDATES(${LANG} "${OpenMPTarget_${LANG}_DEVICE_CANDIDATES}")
+  _OPENMP_TARGET_DEVICE_ARCH_CANDIDATES(${LANG} ${DEVICE} "${OpenMPTarget_${LANG}_DEVICE_CANDIDATES}")
+  _OPENMP_TARGET_WRITE_SOURCE_FILE("${LANG}" "TEST_SOURCE" OpenMPTargetTryFlag _OPENMP_TEST_SRC)
+
+  # Try to compile a test application with the found flags.
+  try_compile(OpenMPTarget_COMPILE_RESULT ${CMAKE_BINARY_DIR} ${_OPENMP_TEST_SRC}
+    CMAKE_FLAGS "-DCOMPILE_DEFINITIONS:STRING=${OpenMP_${LANG}_FLAGS} ${OpenMPTarget_${LANG}_FLAG_CANDIDATES} ${OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES}"
+    "-DINCLUDE_DIRECTORIES:STRING=${OpenMP_${LANG}_INCLUDE_DIR}"
+    LINK_LIBRARIES ${CMAKE_${LANG}_VERBOSE_FLAG}
+    OUTPUT_VARIABLE OpenMP_TRY_COMPILE_OUTPUT
+  )
+
+  file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log
+    "Detecting OpenMP ${CMAKE_${LANG}_COMPILER_ID} ${DEVICE} target support with the following Flags:
+    ${OpenMP_${LANG}_FLAGS} ${OpenMPTarget_${LANG}_FLAG_CANDIDATES} ${OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES}
+    With the following output:\n ${OpenMP_TRY_COMPILE_OUTPUT}\n")
+
+  # If compilation was successful and the device was found set the return variables.
+  if (OpenMPTarget_COMPILE_RESULT AND DEFINED OpenMPTarget_${LANG}_DEVICE_CANDIDATES)
+    file(APPEND ${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/CMakeOutput.log
+      "Compilation successful, adding flags for ${DEVICE}.\n\n")
+
+    # Clang has a seperate library for target offloading.
+    if(CMAKE_${LANG}_COMPILER_ID STREQUAL "Clang")
+      find_library(OpenMPTarget_libomptarget_LIBRARY
+        NAMES omptarget
+        HINTS ${CMAKE_${LANG}_IMPLICIT_LINK_DIRECTORIES}
+      )
+      mark_as_advanced(OpenMPTarget_libomptarget_LIBRARY)
+      set("${OPENMP_LIB_VAR}" "${OpenMPTarget_libomptarget_LIBRARY}" PARENT_SCOPE)
+    else()
+      unset("${OPENMP_LIB_VAR}" PARENT_SCOPE)
+    endif()
+    set("${OPENMP_DEVICE_VAR}" "${OpenMPTarget_${LANG}_DEVICE_CANDIDATES}" PARENT_SCOPE)
+    set("${OPENMP_FLAG_VAR}" "${OpenMPTarget_${LANG}_FLAG_CANDIDATES}" PARENT_SCOPE)
+    set("${OPENMP_ARCH_VAR}" "${OpenMPTarget_${LANG}_DEVICE_ARCH_CANDIDATES}" PARENT_SCOPE)
+  else()
+    unset("${OPENMP_DEVICE_VAR}" PARENT_SCOPE)
+    unset("${OPENMP_FLAG_VAR}" PARENT_SCOPE)
+    unset("${OPENMP_ARCH_VAR}" PARENT_SCOPE)
+  endif()
+endfunction()
+
+# Load the compiler support for each device.
+foreach(LANG IN ITEMS C CXX)
+  # Cache the version in case CMake doesn't load the OpenMP package this time
+  set(OpenMP_${LANG}_VERSION ${OpenMP_${LANG}_VERSION}
+    CACHE STRING "OpenMP Version" FORCE)
+  mark_as_advanced(OpenMP_${LANG}_VERSION)
+  foreach(DEVICE IN ITEMS NVPTX AMDGCN)
+    if(CMAKE_${LANG}_COMPILER_LOADED)
+      if(NOT DEFINED OpenMPTarget_${LANG}_FLAGS OR NOT DEFINED OpenMPTarget_${LANG}_DEVICE)
+        _OPENMP_TARGET_DEVICE_GET_FLAGS(${LANG} ${DEVICE}
+          OpenMPTarget_${DEVICE}_FLAGS_WORK
+          OpenMPTarget_${DEVICE}_LIBS_WORK
+          OpenMPTarget_${DEVICE}_DEVICE_WORK
+          OpenMPTarget_${DEVICE}_ARCHS_WORK)
+
+        separate_arguments(_OpenMPTarget_${DEVICE}_FLAGS NATIVE_COMMAND "${OpenMPTarget_${DEVICE}_FLAGS_WORK}")
+        separate_arguments(_OpenMPTarget_${DEVICE}_ARCHS NATIVE_COMMAND "${OpenMPTarget_${DEVICE}_ARCHS_WORK}")
+        set(OpenMPTarget_${DEVICE}_FLAGS ${_OpenMPTarget_${DEVICE}_FLAGS}
+            CACHE STRING "${DEVICE} target compile flags for OpenMP target offloading" FORCE)
+        set(OpenMPTarget_${DEVICE}_ARCH ${_OpenMPTarget_${DEVICE}_ARCHS}
+            CACHE STRING "${DEVICE} target architecture flags for OpenMP target offloading" FORCE)
+        set(OpenMPTarget_${DEVICE}_LIBS ${OpenMPTarget_${DEVICE}_LIBS_WORK}
+            CACHE STRING "${DEVICE} target libraries for OpenMP target offloading" FORCE)
+        mark_as_advanced(OpenMPTarget_${DEVICE}_FLAGS OpenMPTarget_${DEVICE}_ARCH OpenMPTarget_${DEVICE}_LIBS)
+      endif()
+    endif()
+  endforeach()
+endforeach()
+
+if(OpenMPTarget_FIND_COMPONENTS)
+  set(OpenMPTarget_FINDLIST ${OpenMPTarget_FIND_COMPONENTS})
+else()
+  set(OpenMPTarget_FINDLIST NVPTX)
+endif()
+
+unset(_OpenMPTarget_MIN_VERSION)
+
+# Attempt to find each requested device.
+foreach(LANG IN ITEMS C CXX)
+  foreach(DEVICE IN LISTS OpenMPTarget_FINDLIST)
+    if(CMAKE_${LANG}_COMPILER_LOADED)
+      set(OpenMPTarget_${DEVICE}_VERSION "${OpenMP_${LANG}_VERSION}")
+      set(OpenMPTarget_${DEVICE}_VERSION_MAJOR "${OpenMP_${LANG}_VERSION}_MAJOR")
+      set(OpenMPTarget_${DEVICE}_VERSION_MINOR "${OpenMP_${LANG}_VERSION}_MINOR")
+      set(OpenMPTarget_${DEVICE}_FIND_QUIETLY ${OpenMPTarget_FIND_QUIETLY})
+      set(OpenMPTarget_${DEVICE}_FIND_REQUIRED ${OpenMPTarget_FIND_REQUIRED})
+      set(OpenMPTarget_${DEVICE}_FIND_VERSION ${OpenMPTarget_FIND_VERSION})
+      set(OpenMPTarget_${DEVICE}_FIND_VERSION_EXACT ${OpenMPTarget_FIND_VERSION_EXACT})
+
+      # OpenMP target offloading is only supported in OpenMP 4.0 an newer.
+      if(OpenMPTarget_${DEVICE}_VERSION AND ("${OpenMPTarget_${DEVICE}_VERSION}" VERSION_LESS "4.0"))
+        message(SEND_ERROR "FindOpenMPTarget requires at least OpenMP 4.0")
+      endif()
+
+      set(FPHSA_NAME_MISMATCHED TRUE)
+      find_package_handle_standard_args(OpenMPTarget_${DEVICE}
+        REQUIRED_VARS OpenMPTarget_${DEVICE}_FLAGS
+        VERSION_VAR OpenMPTarget_${DEVICE}_VERSION)
+
+      if(OpenMPTarget_${DEVICE}_FOUND)
+        if(DEFINED OpenMPTarget_${DEVICE}_VERSION)
+          if(NOT _OpenMPTarget_MIN_VERSION OR _OpenMPTarget_MIN_VERSION VERSION_GREATER OpenMPTarget_${LANG}_VERSION)
+            set(_OpenMPTarget_MIN_VERSION OpenMPTarget_${DEVICE}_VERSION)
+          endif()
+        endif()
+        # Create a new target.
+        if(NOT TARGET OpenMPTarget::OpenMPTarget_${DEVICE})
+          add_library(OpenMPTarget::OpenMPTarget_${DEVICE} INTERFACE IMPORTED)
+        endif()
+        # Get compiler flags for offloading to the device and architecture and
+        # set the target features. Include the normal OpenMP flags as well.
+        set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
+          INTERFACE_COMPILE_OPTIONS 
+          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_FLAGS}>"
+          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_ARCH}>"
+          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMP_${LANG}_FLAGS}>")
+        set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
+          INTERFACE_INCLUDE_DIRECTORIES "$<BUILD_INTERFACE:${OpenMP_${LANG}_INCLUDE_DIRS}>")
+        set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
+          INTERFACE_LINK_LIBRARIES 
+          "${OpenMPTarget_${DEVICE}_LIBS}"
+          "${OpenMP_${LANG}_LIBRARIES}")
+        # The offloading flags must also be passed during the linking phase so
+        # the compiler can pass the binary to the correct toolchain.
+        set_property(TARGET OpenMPTarget::OpenMPTarget_${DEVICE} PROPERTY
+          INTERFACE_LINK_OPTIONS 
+          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_FLAGS}>"
+          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMPTarget_${DEVICE}_ARCH}>"
+          "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenMP_${LANG}_FLAGS}>")
+        # Combine all the flags if not using the target for convenience.
+        set(OpenMPTarget_${DEVICE}_FLAGS ${OpenMP_${LANG}_FLAGS}
+          ${OpenMPTarget_${DEVICE}_FLAGS} 
+          ${OpenMPTarget_${DEVICE}_ARCH}
+          CACHE STRING "${DEVICE} target compile flags for OpenMP target offloading" FORCE)
+      endif()
+    endif()
+  endforeach()
+endforeach()
+
+unset(_OpenMPTarget_REQ_VARS)
+foreach(DEVICE IN LISTS OpenMPTarget_FINDLIST)
+  list(APPEND _OpenMPTarget_REQ_VARS "OpenMPTarget_${DEVICE}_FOUND")
+endforeach()
+
+find_package_handle_standard_args(OpenMPTarget
+    REQUIRED_VARS ${_OpenMPTarget_REQ_VARS}
+    VERSION_VAR ${_OpenMPTarget_MIN_VERSION}
+    HANDLE_COMPONENTS)
+
+if(NOT (CMAKE_C_COMPILER_LOADED OR CMAKE_CXX_COMPILER_LOADED) OR CMAKE_Fortran_COMPILER_LOADED)
+  message(SEND_ERROR "FindOpenMPTarget requires the C or CXX languages to be enabled")
+endif()
+
+unset(OpenMPTarget_C_CXX_TEST_SOURCE)
+cmake_policy(POP)

openmp/tools/Modules/README.rst

+=========================
+LLVM OpenMP CMake Modules
+=========================
+
+This directory contains CMake modules for OpenMP. These can be included into a
+project to include different OpenMP features.
+
+.. contents::
+   :local:
+
+Find OpenMP Target Support
+==========================
+
+This module will attempt to find OpenMP target offloading support for a given
+device. The module will attempt to compile a test program using known compiler
+flags for each requested architecture. If successful, the flags required for
+offloading will be loaded into the ``OpenMPTarget::OpenMPTarget_<device>``
+target or the ``OpenMPTarget_NVPTX_FLAGS`` variable. Currently supported target
+devices are ``NVPTX`` and ``AMDGCN``. This module is still under development so
+some features may be missing.
+
+To use this module, simply add the path to CMake's current module path and call
+``find_package``. The module will be installed with your OpenMP installation by
+default. Including OpenMP offloading support in an application should now only
+require a few additions.
+
+.. code-block:: cmake
+
+  cmake_minimum_required(VERSION 3.13.4)
+  project(offloadTest VERSION 1.0 LANGUAGES CXX)
+  
+  list(APPEND CMAKE_MODULE_PATH "${PATH_TO_OPENMP_INSTALL}/lib/cmake/openmp")
+  
+  find_package(OpenMPTarget REQUIRED NVPTX)
+  
+  add_executable(offload)
+  target_link_libraries(offload PRIVATE OpenMPTarget::OpenMPTarget_NVPTX)
+  target_sources(offload PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/src/Main.cpp)
+
+Using this module requires at least CMake version 3.13.4. Supported languages
+are C and C++ with Fortran support planned in the future. If your application
+requires building for a specific device architecture you can set the
+``OpenMPTarget_<device>_ARCH=<flag>`` variable. Compiler support is best for
+Clang but this module should work for other compiler vendors such as IBM or GNU.