Merge topic 'FindOpenMP-offload-target'

14d7dd02d2 FindOpenMP: Add OpenMP::OpenMP_<lang>_Offload targets Acked-by: Kitware Robot <kwrobot@kitware.com> Merge-request: !12718
2026-11-06 05:16:58 +00:00 · 2026-10-30 10:31:51 -04:00
commit a6f8e09e51
@@ -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)`