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

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.

Reviewed By: tianshilei1992

Differential Revision: https://reviews.llvm.org/D104710

3 files changed, 401 insertions, 0 deletions

diff --git a/openmp/tools/Modules/CMakeLists.txt b/openmp/tools/Modules/CMakeLists.txt
new file mode 100644
index 0000000..22d818e
--- /dev/null
+++ b/openmp/tools/Modules/CMakeLists.txt
@@ -0,0 +1,15 @@
+##===----------------------------------------------------------------------===##
+#
+# 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")
diff --git a/openmp/tools/Modules/FindOpenMPTarget.cmake b/openmp/tools/Modules/FindOpenMPTarget.cmake
new file mode 100644
index 0000000..d3917afd
--- /dev/null
+++ b/openmp/tools/Modules/FindOpenMPTarget.cmake
@@ -0,0 +1,342 @@
+##===----------------------------------------------------------------------===##
+#
+# 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)
diff --git a/openmp/tools/Modules/README.rst b/openmp/tools/Modules/README.rst
new file mode 100644
index 0000000..d1ec32f
--- /dev/null
+++ b/openmp/tools/Modules/README.rst
@@ -0,0 +1,44 @@
+=========================
+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.