FindOpenACC: Provide import targets and OpenACC_<lang>_OPTIONS variable

Previously the FindOpenACC module had issues where the contents of
OpenACC_<lang>_FLAGS could not be used with target_compile_options
when it contained multiple compiler flags.

Modules/FindOpenACC.cmake

@@ -22,6 +22,14 @@ project, where ``<lang>`` is one of C, CXX, or Fortran:
   Variable indicating if OpenACC support for ``<lang>`` was detected.
 ``OpenACC_<lang>_FLAGS``
   OpenACC compiler flags for ``<lang>``, separated by spaces.
+``OpenACC_<lang>_OPTIONS``
+  OpenACC compiler flags for ``<lang>``, as a list. Suitable for usage
+  with target_compile_options or target_link_options.
+
+Additionally, the module provides :prop_tgt:`IMPORTED` targets:
+
+``OpenACC::OpenACC_<lang>``
+  Target for using OpenACC from ``<lang>``.
 
 The module will also try to provide the OpenACC version variables:
 
@@ -239,6 +247,9 @@ foreach (LANG IN ITEMS C CXX Fortran)
     if(NOT DEFINED OpenACC_${LANG}_FLAGS)
       _OPENACC_GET_FLAGS("${LANG}" OpenACC_${LANG}_FLAGS)
     endif()
+    if(NOT DEFINED OpenACC_${LANG}_OPTIONS)
+      separate_arguments(OpenACC_${LANG}_OPTIONS NATIVE_COMMAND "${OpenACC_${LANG}_FLAGS}")
+    endif()
     _OPENACC_GET_SPEC_DATE("${LANG}" OpenACC_${LANG}_SPEC_DATE)
     _OPENACC_SET_VERSION_BY_SPEC_DATE("${LANG}")
 
@@ -249,6 +260,23 @@ foreach (LANG IN ITEMS C CXX Fortran)
   endif()
 endforeach()
 
+foreach (LANG IN ITEMS C CXX Fortran)
+  if(OpenACC_${LANG}_FOUND AND NOT TARGET OpenACC::OpenACC_${LANG})
+    add_library(OpenACC::OpenACC_${LANG} INTERFACE IMPORTED)
+  endif()
+  if(OpenACC_${LANG}_LIBRARIES)
+    set_property(TARGET OpenACC::OpenACC_${LANG} PROPERTY
+      INTERFACE_LINK_LIBRARIES "${OpenACC_${LANG}_LIBRARIES}")
+  endif()
+  if(OpenACC_${LANG}_FLAGS)
+    set_property(TARGET OpenACC::OpenACC_${LANG} PROPERTY
+      INTERFACE_COMPILE_OPTIONS "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenACC_${LANG}_OPTIONS}>")
+    set_property(TARGET OpenACC::OpenACC_${LANG} PROPERTY
+      INTERFACE_LINK_OPTIONS "$<$<COMPILE_LANGUAGE:${LANG}>:${OpenACC_${LANG}_OPTIONS}>")
+    unset(_OpenACC_${LANG}_OPTIONS)
+  endif()
+endforeach()
+
 unset(OpenACC_C_CXX_TEST_SOURCE)
 unset(OpenACC_Fortran_TEST_SOURCE)
 unset(OpenACC_C_CXX_CHECK_VERSION_SOURCE)

Tests/CMakeLists.txt

@@ -1409,6 +1409,7 @@ ${CMake_SOURCE_DIR}/Utilities/Release/push.bash --dir dev -- '${CMake_BUILD_NIGH
             LibXml2
             LTTngUST
             ODBC
+            OpenACC
             OpenCL
             OpenGL
             OpenMP

Tests/FindOpenACC/CMakeLists.txt

@@ -0,0 +1,20 @@
+
+set(langs C CXX)
+if(NOT CMAKE_GENERATOR STREQUAL "Ninja")
+  list(APPEND langs Fortran)
+endif()
+
+foreach(lang IN LISTS langs)
+  if(CMAKE_${lang}_COMPILER)
+    add_test(NAME FindOpenACC.Test${lang} COMMAND
+      ${CMAKE_CTEST_COMMAND} -C $<CONFIGURATION>
+      --build-and-test
+      "${CMake_SOURCE_DIR}/Tests/FindOpenACC/${lang}Test"
+      "${CMake_BINARY_DIR}/Tests/FindOpenACC/${lang}Test"
+      ${build_generator_args}
+      --build-project TestFindOpenACC
+      --build-options ${build_options}
+      --test-command ${CMAKE_CTEST_COMMAND} -V -C $<CONFIGURATION>
+      )
+  endif()
+endforeach()

Tests/FindOpenACC/CTest/CMakeLists.txt

@@ -0,0 +1,13 @@
+cmake_minimum_required(VERSION 3.14)
+project(VerifyFindOpenAcc C)
+
+set(CMAKE_C_STANDARD 11)
+
+find_package(OpenACC REQUIRED)
+
+add_executable(UsesOpenACC main.c)
+target_link_libraries(UsesOpenACC PRIVATE OpenACC::OpenACC_C)
+
+add_executable(UsesOpenACCVars main.c)
+target_link_options(UsesOpenACCVars PRIVATE ${OpenACC_C_OPTIONS})
+target_compile_options(UsesOpenACCVars PRIVATE ${OpenACC_C_OPTIONS})

Tests/FindOpenACC/CTest/main.c

@@ -0,0 +1,44 @@
+#include <stdio.h>
+#include <stdlib.h>
+
+void vecaddgpu(float* r, float* a, float* b, int n)
+{
+#pragma acc kernels loop present(r, a, b)
+  for (int i = 0; i < n; ++i)
+    r[i] = a[i] + b[i];
+}
+
+int main()
+{
+  int n = 100000; /* vector length */
+  float* a;       /* input vector 1 */
+  float* b;       /* input vector 2 */
+  float* r;       /* output vector */
+  float* e;       /* expected output values */
+  int i, errs;
+
+  a = (float*)malloc(n * sizeof(float));
+  b = (float*)malloc(n * sizeof(float));
+  r = (float*)malloc(n * sizeof(float));
+  e = (float*)malloc(n * sizeof(float));
+  for (i = 0; i < n; ++i) {
+    a[i] = (float)(i + 1);
+    b[i] = (float)(1000 * i);
+  }
+/* compute on the GPU */
+#pragma acc data copyin(a [0:n], b [0:n]) copyout(r [0:n])
+  {
+    vecaddgpu(r, a, b, n);
+  }
+  /* compute on the host to compare */
+  for (i = 0; i < n; ++i)
+    e[i] = a[i] + b[i];
+  /* compare results */
+  errs = 0;
+  for (i = 0; i < n; ++i) {
+    if (r[i] != e[i]) {
+      ++errs;
+    }
+  }
+  return errs;
+}

Tests/FindOpenACC/CXXTest/CMakeLists.txt

@@ -0,0 +1,13 @@
+cmake_minimum_required(VERSION 3.14)
+project(VerifyFindOpenAcc CXX)
+
+set(CMAKE_CXX_STANDARD 11)
+
+find_package(OpenACC REQUIRED)
+
+add_executable(UsesOpenACC main.cxx)
+target_link_libraries(UsesOpenACC PRIVATE OpenACC::OpenACC_CXX)
+
+add_executable(UsesOpenACCVars main.cxx)
+target_link_options(UsesOpenACCVars PRIVATE ${OpenACC_CXX_OPTIONS})
+target_compile_options(UsesOpenACCVars PRIVATE ${OpenACC_CXX_OPTIONS})

Tests/FindOpenACC/CXXTest/main.cxx

@@ -0,0 +1,43 @@
+
+#include <vector>
+
+void vecaddgpu(float* r, float* a, float* b, std::size_t n)
+{
+#pragma acc kernels loop present(r, a, b)
+  for (std::size_t i = 0; i < n; ++i)
+    r[i] = a[i] + b[i];
+}
+
+int main(int, char* [])
+{
+  const std::size_t n = 100000; /* vector length */
+  std::vector<float> a(n);      /* input vector 1 */
+  std::vector<float> b(n);      /* input vector 2 */
+  std::vector<float> r(n);      /* output vector */
+  std::vector<float> e(n);      /* expected output values */
+
+  for (std::size_t i = 0; i < n; ++i) {
+    a[i] = static_cast<float>(i + 1);
+    b[i] = static_cast<float>(1000 * i);
+  }
+
+  /* compute on the GPU */
+  auto a_ptr = a.data();
+  auto b_ptr = b.data();
+  auto r_ptr = r.data();
+#pragma acc data copyin(a_ptr [0:n], b_ptr [0:n]) copyout(r_ptr [0:n])
+  {
+    vecaddgpu(r_ptr, a_ptr, b_ptr, n);
+  }
+  /* compute on the host to compare */
+  for (std::size_t i = 0; i < n; ++i)
+    e[i] = a[i] + b[i];
+  /* compare results */
+  int errs = 0;
+  for (std::size_t i = 0; i < n; ++i) {
+    if (r[i] != e[i]) {
+      ++errs;
+    }
+  }
+  return errs;
+}

Tests/FindOpenACC/FortranTest/CMakeLists.txt

@@ -0,0 +1,11 @@
+cmake_minimum_required(VERSION 3.14)
+project(VerifyFindOpenAcc Fortran)
+
+find_package(OpenACC REQUIRED)
+
+add_executable(UsesOpenACC main.f90)
+target_link_libraries(UsesOpenACC PRIVATE OpenACC::OpenACC_Fortran)
+
+add_executable(UsesOpenACCVars main.f90)
+target_link_options(UsesOpenACCVars PRIVATE ${OpenACC_Fortran_OPTIONS})
+target_compile_options(UsesOpenACCVars PRIVATE ${OpenACC_Fortran_OPTIONS})

Tests/FindOpenACC/FortranTest/main.f90

@@ -0,0 +1,9 @@
+program t
+integer(4) a(10000)
+a = [ (1+i,i=1,10000) ]
+!$acc kernels
+do i = 1, 10000
+  if (a(i)/3000*3000.eq.a(i)) print *," located ",i,a(i),i.gt.5000,a(i)/5.0
+end do
+!$acc end kernels
+end