FindOpenMP: Add OpenMP::OpenMP_<lang>_Offload targets

OpenMP::OpenMP_<lang> carries no device flags, and for GCC and Clang links the OpenMP runtime without the OpenMP flag, so target regions built through it fail to link or silently run on the host. Once OpenMP_<lang>_OFFLOAD or OpenMP_<lang>_OFFLOAD_FLAGS is set, the new targets add the device flags, plus the OpenMP flag at link time; otherwise they match the host targets. FindOpenMP.Offload runs a target kernel on a GPU and is enabled in the nvhpc job. Fixes: #21784, #26701
2026-11-06 05:16:58 +00:00 · 2026-10-29 17:12:29 +01:00
1 parent 6658c5e419
commit 14d7dd02d2
@@ -5,6 +5,14 @@ set(CMake_TEST_CUDA_CUPTI "ON" CACHE STRING "")
 set(CMake_TEST_C_STANDARDS "90;99;11;17;23" CACHE STRING "")
 set(CMake_TEST_CXX_STANDARDS "98;11;14;17;20;23;26" CACHE STRING "")

+set(CMake_TEST_FindOpenMP "ON" CACHE BOOL "")
+set(CMake_TEST_FindOpenMP_C "ON" CACHE BOOL "")
+set(CMake_TEST_FindOpenMP_CXX "ON" CACHE BOOL "")
+set(CMake_TEST_FindOpenMP_Fortran "ON" CACHE BOOL "")
+set(CMake_TEST_FindOpenMP_C_OFFLOAD "ON" CACHE BOOL "")
+set(CMake_TEST_FindOpenMP_CXX_OFFLOAD "ON" CACHE BOOL "")
+set(CMake_TEST_FindOpenMP_Fortran_OFFLOAD "ON" CACHE BOOL "")
+
 set(configure_no_sccache 1)

 include("${CMAKE_CURRENT_LIST_DIR}/configure_external_test.cmake")
@@ -0,0 +1,7 @@
+FindOpenMP-offload-target
+-------------------------
+
+* The :module:`FindOpenMP` module now provides ``OpenMP::OpenMP_<lang>_Offload``
+  imported targets for code whose OpenMP ``target`` regions may be offloaded
+  to a device.  Offload is enabled with the new ``OpenMP_<lang>_OFFLOAD`` and
+  ``OpenMP_<lang>_OFFLOAD_FLAGS`` variables.
@@ -59,6 +59,27 @@ This module provides the following :ref:`Imported Targets`:

  Target encapsulating the OpenMP usage requirements for language ``<lang>``.

+``OpenMP::OpenMP_<lang>_Offload``
+  .. versionadded:: 4.6
+
+  Target for code whose OpenMP ``target`` regions may be offloaded to a
+  device, for ``<lang>`` of ``C``, ``CXX``, or ``Fortran``.  It links
+  ``OpenMP::OpenMP_<lang>`` and, when offload is enabled with the
+  ``OpenMP_<lang>_OFFLOAD`` or ``OpenMP_<lang>_OFFLOAD_FLAGS`` variables,
+  adds the offload flags when compiling ``<lang>`` sources.  It also adds
+  them, together with the OpenMP flags, when the link is done by a compiler
+  with the same :variable:`CMAKE_<LANG>_COMPILER_ID`, because compilers
+  need both in order to link device code.  When offload is not enabled, it
+  is equivalent to ``OpenMP::OpenMP_<lang>``.
+
+  With NVHPC, an executable linked by NVHPC must also link the offload
+  target when it uses a shared library with offloaded code, for example
+  through ``$<LINK_ONLY:OpenMP::OpenMP_<lang>_Offload>`` in the library's
+  ``INTERFACE`` link libraries.
+
+  Linking through a separate linker, as with Visual Studio generators or
+  ``clang-cl``, is not supported.
+
 Input Variables
 ^^^^^^^^^^^^^^^

@@ -75,6 +96,49 @@ The following variables may be set to influence this module's behavior:

 .. _`msvc-openmp`: https://learn.microsoft.com/en-us/cpp/build/reference/openmp-enable-openmp-2-0-support

+``OpenMP_<lang>_OFFLOAD``
+  .. versionadded:: 4.6
+
+  Enables OpenMP ``target`` offload for targets that link
+  ``OpenMP::OpenMP_<lang>_Offload``.  If not defined, offload is enabled
+  when ``OpenMP_<lang>_OFFLOAD_FLAGS`` is non-empty.  If defined, its
+  boolean value decides, so ``OFF`` disables offload even when flags are
+  set.
+
+  The offload flags are ``OpenMP_<lang>_OFFLOAD_FLAGS`` if non-empty, and
+  otherwise a default for the compiler:
+
+  * GNU: none.  GCC offloads to the devices whose offload compilers are
+    installed, such as ``gcc-offload-nvptx`` on Debian and Ubuntu.
+  * NVHPC: ``-mp=gpu``.
+  * IntelLLVM with a GNU-like command line (``icx``, ``icpx``, ``ifx``):
+    ``-fopenmp-targets=spir64``.
+  * Cray, CrayClang: none.  The device comes from the loaded
+    ``craype-accel-*`` module.
+
+  Other compilers, including Clang, have no default.  Enabling offload for
+  them without ``OpenMP_<lang>_OFFLOAD_FLAGS`` is an error, even if the
+  package is ``QUIET`` or the component is optional.
+
+  Targets linking only ``OpenMP::OpenMP_<lang>`` are not affected.
+  Projects using the result variables must add the device flags for
+  compiling and linking themselves.  Disabling offload only removes what
+  the offload target would add; it does not undo offload enabled by other
+  compiler flags or by the environment.  ``OpenMP_<lang>_FOUND`` and
+  ``OpenMP_<lang>_VERSION`` describe host OpenMP support only; the module
+  does not check that offload works.
+
+``OpenMP_<lang>_OFFLOAD_FLAGS``
+  .. versionadded:: 4.6
+
+  Compiler flags selecting the device(s) to offload OpenMP ``target``
+  regions to, replacing the compiler's default described above.  For
+  example, ``--offload-arch=sm_80`` for Clang, ``-foffload=nvptx-none`` for
+  GCC, or ``-mp=gpu -gpu=cc80`` for NVHPC.  The value is a command-line
+  string, split into arguments as by the ``SHELL:`` prefix of
+  :command:`target_compile_options`.  An empty or whitespace-only value
+  counts as not set.
+
 Result Variables
 ^^^^^^^^^^^^^^^^

@@ -203,6 +267,19 @@ using the C language component:

  find_package(OpenMP)
  target_link_libraries(project_target PRIVATE OpenMP::OpenMP_C)
+
+Linking a library whose ``target`` regions may be offloaded to a GPU:
+
+.. code-block:: cmake
+
+  find_package(OpenMP REQUIRED)
+  target_link_libraries(kernels PRIVATE OpenMP::OpenMP_CXX_Offload)
+
+and enabling offload when configuring the build, e.g. with Clang:
+
+.. code-block:: console
+
+  $ cmake -B build -DOpenMP_CXX_OFFLOAD_FLAGS=--offload-arch=sm_80
 #]=======================================================================]

 cmake_policy(PUSH)
@@ -824,6 +901,69 @@ foreach(LANG IN LISTS OpenMP_FINDLIST)
        set_property(TARGET OpenMP::OpenMP_${LANG} PROPERTY
          INTERFACE_LINK_LIBRARIES "${OpenMP_${LANG}_LIBRARIES}")
      endif()
+
+      if(NOT LANG STREQUAL "CUDA")
+        if(NOT TARGET OpenMP::OpenMP_${LANG}_Offload)
+          add_library(OpenMP::OpenMP_${LANG}_Offload INTERFACE IMPORTED)
+          set_property(TARGET OpenMP::OpenMP_${LANG}_Offload PROPERTY
+            INTERFACE_LINK_LIBRARIES OpenMP::OpenMP_${LANG})
+        endif()
+        string(STRIP "${OpenMP_${LANG}_OFFLOAD_FLAGS}" _OpenMP_OFFLOAD_FLAGS)
+        if(DEFINED OpenMP_${LANG}_OFFLOAD)
+          set(_OpenMP_OFFLOAD "${OpenMP_${LANG}_OFFLOAD}")
+        elseif(NOT _OpenMP_OFFLOAD_FLAGS STREQUAL "")
+          set(_OpenMP_OFFLOAD ON)
+        else()
+          set(_OpenMP_OFFLOAD OFF)
+        endif()
+        # A repeated find_package() with different inputs replaces these
+        # options rather than adding to them.
+        set(_OpenMP_COMPILE_OPTIONS "")
+        set(_OpenMP_LINK_OPTIONS "")
+        if(_OpenMP_OFFLOAD)
+          # GCC and Cray need no default: GCC offloads to whichever offload
+          # compilers are installed, and Cray's wrappers take the device from
+          # the craype-accel module.
+          if(_OpenMP_OFFLOAD_FLAGS STREQUAL ""
+              AND NOT CMAKE_${LANG}_COMPILER_ID MATCHES "^(GNU|Cray|CrayClang)$")
+            if(CMAKE_${LANG}_COMPILER_ID STREQUAL "NVHPC")
+              set(_OpenMP_OFFLOAD_FLAGS "-mp=gpu")
+            elseif(CMAKE_${LANG}_COMPILER_ID STREQUAL "IntelLLVM"
+                AND NOT "x${CMAKE_${LANG}_COMPILER_FRONTEND_VARIANT}" STREQUAL "xMSVC")
+              set(_OpenMP_OFFLOAD_FLAGS "-fopenmp-targets=spir64")
+            else()
+              message(FATAL_ERROR
+                "OpenMP_${LANG}_OFFLOAD is enabled, but FindOpenMP has no default "
+                "offload flags for the ${CMAKE_${LANG}_COMPILER_ID} compiler.  "
+                "Set OpenMP_${LANG}_OFFLOAD_FLAGS to select the device, e.g. "
+                "--offload-arch=<arch> for Clang.")
+            endif()
+          endif()
+          # Links by the same compiler family, e.g. a C kernel library in a
+          # C++ executable, take these flags; another vendor's driver can't.
+          set(_OpenMP_LINK_DRIVER "$<OR:$<LINK_LANG_AND_ID:C,${CMAKE_${LANG}_COMPILER_ID}>,$<LINK_LANG_AND_ID:CXX,${CMAKE_${LANG}_COMPILER_ID}>,$<LINK_LANG_AND_ID:Fortran,${CMAKE_${LANG}_COMPILER_ID}>>")
+          # The compiler driver only runs the offload link step if it sees
+          # the OpenMP flags when linking.
+          string(STRIP "${OpenMP_${LANG}_FLAGS}" _OpenMP_FLAGS)
+          if(NOT _OpenMP_FLAGS STREQUAL "")
+            list(APPEND _OpenMP_LINK_OPTIONS "$<${_OpenMP_LINK_DRIVER}:SHELL:${_OpenMP_FLAGS}>")
+          endif()
+          if(NOT _OpenMP_OFFLOAD_FLAGS STREQUAL "")
+            list(APPEND _OpenMP_COMPILE_OPTIONS "$<$<COMPILE_LANGUAGE:${LANG}>:SHELL:${_OpenMP_OFFLOAD_FLAGS}>")
+            list(APPEND _OpenMP_LINK_OPTIONS "$<${_OpenMP_LINK_DRIVER}:SHELL:${_OpenMP_OFFLOAD_FLAGS}>")
+          endif()
+          unset(_OpenMP_FLAGS)
+          unset(_OpenMP_LINK_DRIVER)
+        endif()
+        set_property(TARGET OpenMP::OpenMP_${LANG}_Offload PROPERTY
+          INTERFACE_COMPILE_OPTIONS "${_OpenMP_COMPILE_OPTIONS}")
+        set_property(TARGET OpenMP::OpenMP_${LANG}_Offload PROPERTY
+          INTERFACE_LINK_OPTIONS "${_OpenMP_LINK_OPTIONS}")
+        unset(_OpenMP_COMPILE_OPTIONS)
+        unset(_OpenMP_LINK_OPTIONS)
+        unset(_OpenMP_OFFLOAD_FLAGS)
+        unset(_OpenMP_OFFLOAD)
+      endif()
    endif()
  endif()
 endforeach()
@@ -26,3 +26,26 @@ endif()
 if(CMake_TEST_FindOpenMP_FLAG_Fortran)
  set_property(TEST FindOpenMP.Test APPEND PROPERTY LABELS "Fortran")
 endif()
+
+unset(offload_options)
+foreach(c C CXX Fortran)
+  if(CMake_TEST_FindOpenMP_${c}_OFFLOAD)
+    list(APPEND offload_options -DOpenMP_TEST_OFFLOAD_${c}=1)
+    if(CMake_TEST_FindOpenMP_${c}_OFFLOAD_FLAGS)
+      list(APPEND offload_options "-DOpenMP_${c}_OFFLOAD_FLAGS=${CMake_TEST_FindOpenMP_${c}_OFFLOAD_FLAGS}")
+    endif()
+  endif()
+endforeach()
+if(offload_options)
+  add_test(NAME FindOpenMP.Offload COMMAND
+    ${CMAKE_CTEST_COMMAND} -C $<CONFIGURATION>
+    --build-and-test
+    "${CMake_SOURCE_DIR}/Tests/FindOpenMP/Offload"
+    "${CMake_BINARY_DIR}/Tests/FindOpenMP/Offload"
+    ${build_generator_args}
+    --build-project TestFindOpenMPOffload
+    --build-options ${build_options} ${offload_options}
+    --test-command ${CMAKE_CTEST_COMMAND} -V -C $<CONFIGURATION>
+    )
+  set_property(TEST FindOpenMP.Offload APPEND PROPERTY LABELS "OpenMPOffload")
+endif()
@@ -0,0 +1,79 @@
+cmake_minimum_required(VERSION 4.0)
+project(TestFindOpenMPOffload NONE)
+include(CTest)
+
+foreach(c C CXX Fortran)
+  if(OpenMP_TEST_OFFLOAD_${c})
+    enable_language(${c})
+    list(APPEND langs ${c})
+    # Offload is enabled here as a builder would; the targets below opt in
+    # by linking the offload target.
+    set(OpenMP_${c}_OFFLOAD ON)
+  endif()
+endforeach()
+
+find_package(OpenMP REQUIRED COMPONENTS ${langs})
+
+set(main_C main.c)
+set(offload_C offload.c)
+if("CXX" IN_LIST langs)
+  configure_file(main.c main.cxx COPYONLY)
+  configure_file(offload.c offload.cxx COPYONLY)
+  set(main_CXX ${CMAKE_CURRENT_BINARY_DIR}/main.cxx)
+  set(offload_CXX ${CMAKE_CURRENT_BINARY_DIR}/offload.cxx)
+endif()
+if("Fortran" IN_LIST langs)
+  if(OpenMP_Fortran_HAVE_OMPLIB_MODULE)
+    set(OpenMP_Fortran_INCLUDE_LINE "use omp_lib\n      implicit none")
+  else()
+    set(OpenMP_Fortran_INCLUDE_LINE "implicit none\n      include 'omp_lib.h'")
+  endif()
+  configure_file(offload.f90.in offload.f90 @ONLY)
+  set(main_Fortran main.f90)
+  set(offload_Fortran ${CMAKE_CURRENT_BINARY_DIR}/offload.f90)
+endif()
+
+# Each compiler handles the offload flags for its own language.
+foreach(c IN LISTS langs)
+  add_executable(offload_exe_${c} ${main_${c}} ${offload_${c}})
+  target_link_libraries(offload_exe_${c} PRIVATE OpenMP::OpenMP_${c}_Offload)
+  list(APPEND tests exe_${c})
+endforeach()
+
+# How the flags reach the link that embeds the device code doesn't depend
+# on the language, so the library layouts are only built for C.
+if("C" IN_LIST langs)
+  # Static library: the flags reach the executable's link through the
+  # library's usage requirements.
+  add_library(offload_lib STATIC offload.c)
+  target_link_libraries(offload_lib PRIVATE OpenMP::OpenMP_C_Offload)
+  add_executable(offload_static main.c)
+  target_link_libraries(offload_static PRIVATE offload_lib)
+
+  # Shared library: its own link embeds the device code.
+  add_library(offload_dso SHARED offload.c)
+  target_link_libraries(offload_dso PRIVATE OpenMP::OpenMP_C_Offload)
+  add_executable(offload_shared main.c)
+  target_link_libraries(offload_shared PRIVATE offload_dso)
+  if(CMAKE_C_COMPILER_ID STREQUAL "NVHPC")
+    # Without -mp=gpu, NVHPC links host-only stubs of its GPU runtime into
+    # the executable, and they shadow the ones the library needs.
+    target_link_libraries(offload_shared PRIVATE
+      "$<LINK_ONLY:OpenMP::OpenMP_C_Offload>")
+  endif()
+
+  list(APPEND tests static shared)
+
+  if("CXX" IN_LIST langs AND CMAKE_C_COMPILER_ID STREQUAL CMAKE_CXX_COMPILER_ID)
+    # C kernel library linked by the C++ compiler driver.
+    add_executable(offload_mixed ${main_CXX})
+    target_link_libraries(offload_mixed PRIVATE offload_lib)
+    list(APPEND tests mixed)
+  endif()
+endif()
+
+foreach(t IN LISTS tests)
+  add_test(NAME offload_${t} COMMAND offload_${t})
+  # A silent fallback to the host fails the test instead of passing it.
+  set_property(TEST offload_${t} PROPERTY ENVIRONMENT OMP_TARGET_OFFLOAD=MANDATORY)
+endforeach()
@@ -0,0 +1,10 @@
+#ifdef __cplusplus
+extern "C"
+#endif
+  int
+  offload_check(void);
+
+int main(void)
+{
+  return offload_check();
+}
@@ -0,0 +1,5 @@
+      program main
+      implicit none
+      integer, external :: offload_check
+      if (offload_check() /= 0) error stop 1
+      end program main
@@ -0,0 +1,20 @@
+#include <omp.h>
+
+#ifdef __cplusplus
+extern "C"
+#endif
+  int
+  offload_check(void)
+{
+  int i;
+  int on_host = 1;
+  int sum = 0;
+#pragma omp target map(tofrom : on_host, sum)
+  {
+    on_host = omp_is_initial_device();
+    for (i = 1; i <= 10; ++i) {
+      sum += i;
+    }
+  }
+  return on_host || sum != 55;
+}
@@ -0,0 +1,15 @@
+      integer function offload_check()
+      @OpenMP_Fortran_INCLUDE_LINE@
+      integer :: i, s
+      logical :: on_host
+      on_host = .true.
+      s = 0
+!$omp target map(tofrom: on_host, s)
+      on_host = omp_is_initial_device()
+      do i = 1, 10
+        s = s + i
+      end do
+!$omp end target
+      offload_check = 0
+      if (on_host .or. s /= 55) offload_check = 1
+      end function offload_check
@@ -54,7 +54,7 @@ if(CMAKE_C_COMPILER_ID STREQUAL "Clang" AND CMAKE_C_SIMULATE_ID STREQUAL "MSVC")
  endif()
 endif()

-# FIXME(#26701) Needed to check target offload conditions
+# Cray needs the OpenMP flags at link time to use its offload linker
 foreach(c C CXX Fortran)
  if(NOT "${OpenMP_TEST_${c}}")
    continue()
@@ -558,6 +558,11 @@ if(CMake_TEST_FindRuby)
    )
 endif()
 add_RunCMake_test(FindOpenGL)
+if(CMake_TEST_FindOpenMP_C AND CMake_TEST_FindOpenMP_CXX
+    AND CMAKE_C_COMPILER_ID MATCHES "^(GNU|Clang|AppleClang)$"
+    AND NOT WIN32 AND CMAKE_GENERATOR MATCHES "Make|Ninja")
+  add_RunCMake_test(FindOpenMP -DCMAKE_C_COMPILER_ID=${CMAKE_C_COMPILER_ID})
+endif()
 add_RunCMake_test(FindTclStub)
 add_RunCMake_test(InitialFlags)
 if(CMake_TEST_FindOpenSSL)
@@ -0,0 +1,3 @@
+cmake_minimum_required(VERSION 4.0)
+project(${RunCMake_TEST} C CXX)
+include(${RunCMake_TEST}.cmake)
@@ -0,0 +1,5 @@
+set(mark MARK_A)
+set(link_omp exe static mixed)
+set(link_mark exe static mixed)
+set(compile_mark exe lib)
+include(${CMAKE_CURRENT_LIST_DIR}/check-flags.cmake)
@@ -0,0 +1,2 @@
+set(OpenMP_C_OFFLOAD_FLAGS -DMARK_A)
+include(${CMAKE_CURRENT_LIST_DIR}/Offload-common.cmake)
@@ -0,0 +1 @@
+1
@@ -0,0 +1,4 @@
+^CMake Error at [^
+]*/FindOpenMP\.cmake:[0-9]+ \(message\):
+  OpenMP_C_OFFLOAD is enabled, but FindOpenMP has no default offload flags
+  for the [A-Za-z]+ compiler\.
@@ -0,0 +1,2 @@
+set(OpenMP_C_OFFLOAD ON)
+include(${CMAKE_CURRENT_LIST_DIR}/Offload-common.cmake)
@@ -0,0 +1 @@
+include(${CMAKE_CURRENT_LIST_DIR}/check-flags.cmake)
@@ -0,0 +1,3 @@
+set(OpenMP_C_OFFLOAD OFF)
+set(OpenMP_C_OFFLOAD_FLAGS -DMARK_A)
+include(${CMAKE_CURRENT_LIST_DIR}/Offload-common.cmake)
@@ -0,0 +1 @@
+include(${CMAKE_CURRENT_LIST_DIR}/check-flags.cmake)
@@ -0,0 +1,3 @@
+set(OpenMP_C_OFFLOAD "")
+set(OpenMP_C_OFFLOAD_FLAGS -DMARK_A)
+include(${CMAKE_CURRENT_LIST_DIR}/Offload-common.cmake)
@@ -0,0 +1,17 @@
+find_package(OpenMP REQUIRED COMPONENTS C CXX)
+
+add_executable(exe main.c)
+target_link_libraries(exe PRIVATE OpenMP::OpenMP_C_Offload)
+
+add_library(lib STATIC lib.c)
+target_link_libraries(lib PRIVATE OpenMP::OpenMP_C_Offload)
+add_executable(static main.c)
+target_link_libraries(static PRIVATE lib)
+add_executable(mixed main.cxx)
+target_link_libraries(mixed PRIVATE lib)
+
+# Not opted in, so these get no offload flags.
+add_executable(hostc main.c)
+target_link_libraries(hostc PRIVATE OpenMP::OpenMP_C)
+add_executable(host main.cxx)
+target_link_libraries(host PRIVATE OpenMP::OpenMP_CXX)
@@ -0,0 +1,5 @@
+set(mark MARK_B)
+set(link_omp exe static mixed)
+set(link_mark exe static mixed)
+set(compile_mark exe lib)
+include(${CMAKE_CURRENT_LIST_DIR}/check-flags.cmake)
@@ -0,0 +1,5 @@
+set(OpenMP_C_OFFLOAD_FLAGS -DMARK_A)
+include(${CMAKE_CURRENT_LIST_DIR}/Offload-common.cmake)
+
+set(OpenMP_C_OFFLOAD_FLAGS -DMARK_B)
+find_package(OpenMP REQUIRED COMPONENTS C CXX)
@@ -0,0 +1 @@
+include(${CMAKE_CURRENT_LIST_DIR}/check-flags.cmake)
@@ -0,0 +1,5 @@
+set(OpenMP_C_OFFLOAD_FLAGS -DMARK_A)
+include(${CMAKE_CURRENT_LIST_DIR}/Offload-common.cmake)
+
+set(OpenMP_C_OFFLOAD OFF)
+find_package(OpenMP REQUIRED COMPONENTS C CXX)
@@ -0,0 +1,21 @@
+include(RunCMake)
+
+function(run_offload_build case)
+  run_cmake(${case})
+  set(RunCMake_TEST_BINARY_DIR ${RunCMake_BINARY_DIR}/${case}-build)
+  set(RunCMake_TEST_NO_CLEAN 1)
+  run_cmake_command(${case}-build ${CMAKE_COMMAND} --build . --verbose)
+endfunction()
+
+run_offload_build(FlagsImplyOn)
+if(CMAKE_C_COMPILER_ID STREQUAL "GNU")
+  run_offload_build(SwitchOn)
+else()
+  run_cmake(NoDefault)
+endif()
+run_offload_build(Off)
+run_offload_build(OffEmpty)
+run_offload_build(Unset)
+run_offload_build(WhitespaceFlags)
+run_offload_build(RefindFlags)
+run_offload_build(RefindOff)
@@ -0,0 +1,2 @@
+set(link_omp exe static mixed)
+include(${CMAKE_CURRENT_LIST_DIR}/check-flags.cmake)
@@ -0,0 +1,2 @@
+set(OpenMP_C_OFFLOAD ON)
+include(${CMAKE_CURRENT_LIST_DIR}/Offload-common.cmake)
@@ -0,0 +1 @@
+include(${CMAKE_CURRENT_LIST_DIR}/check-flags.cmake)
@@ -0,0 +1 @@
+include(${CMAKE_CURRENT_LIST_DIR}/Offload-common.cmake)
@@ -0,0 +1 @@
+include(${CMAKE_CURRENT_LIST_DIR}/check-flags.cmake)
@@ -0,0 +1,2 @@
+set(OpenMP_C_OFFLOAD_FLAGS "   ")
+include(${CMAKE_CURRENT_LIST_DIR}/Offload-common.cmake)
@@ -0,0 +1,60 @@
+# Checks the verbose build output against these lists of targets:
+#   link_omp     - link line has the OpenMP flag
+#   link_mark    - link line has ${mark}
+#   compile_mark - compile line has ${mark}
+# Other targets lack them, and no other marker appears.
+
+string(REPLACE ";" "," out "${actual_stdout}")
+string(REPLACE "\n" ";" lines "${out}")
+
+function(check_line kind t line what pattern want)
+  list(FIND ${want} "${t}" i)
+  if(line MATCHES "${pattern}")
+    if(i EQUAL -1)
+      string(APPEND msg "${kind} line of '${t}' has ${what} but should not:\n  ${line}\n")
+    endif()
+  elseif(NOT i EQUAL -1)
+    string(APPEND msg "${kind} line of '${t}' lacks ${what}:\n  ${line}\n")
+  endif()
+  set(msg "${msg}" PARENT_SCOPE)
+endfunction()
+
+set(msg "")
+foreach(t IN ITEMS exe lib static mixed hostc host)
+  set(compile "")
+  set(link "")
+  foreach(l IN LISTS lines)
+    if(l MATCHES " -o CMakeFiles/${t}\\.dir/[^ ]+ -c ")
+      set(compile "${l}")
+    elseif(l MATCHES " -o ([^ ]*/)?${t}( |$)")
+      set(link "${l}")
+    endif()
+  endforeach()
+  if(compile STREQUAL "")
+    string(APPEND msg "No compile line for '${t}'.\n")
+  elseif(mark)
+    check_line(Compile ${t} "${compile}" "${mark}" "-D${mark}( |$)" compile_mark)
+  endif()
+  # The static library is archived, not linked.
+  if(t STREQUAL "lib")
+    continue()
+  endif()
+  if(link STREQUAL "")
+    string(APPEND msg "No link line for '${t}'.\n")
+    continue()
+  endif()
+  check_line(Link ${t} "${link}" "the OpenMP flag" "fopenmp" link_omp)
+  if(mark)
+    check_line(Link ${t} "${link}" "${mark}" "-D${mark}( |$)" link_mark)
+  endif()
+endforeach()
+
+foreach(m IN ITEMS MARK_A MARK_B)
+  if(NOT m STREQUAL "${mark}" AND actual_stdout MATCHES "-D${m}")
+    string(APPEND msg "Unexpected ${m} in the build output.\n")
+  endif()
+endforeach()
+
+if(msg)
+  set(RunCMake_TEST_FAILED "${msg}")
+endif()
@@ -0,0 +1,4 @@
+int lib(void)
+{
+  return 0;
+}
@@ -0,0 +1,4 @@
+int main(void)
+{
+  return 0;
+}
@@ -0,0 +1,4 @@
+int main()
+{
+  return 0;
+}
				`@@ -0,0 +1 @@`
				`include(${CMAKE_CURRENT_LIST_DIR}/check-flags.cmake)`
				`@@ -0,0 +1 @@`
				`include(${CMAKE_CURRENT_LIST_DIR}/Offload-common.cmake)`