OpenMP, C++: Add template support for the has_device_addr clause.

This patch adds support for list items in the has_device_addr clause which type
is given by C++ template parameters.

gcc/cp/ChangeLog:

	* pt.cc (tsubst_omp_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR.
	* semantics.cc (finish_omp_clauses): Added template decl processing.

libgomp/ChangeLog:

	* testsuite/libgomp.c++/target-has-device-addr-7.C: New test.
	* testsuite/libgomp.c++/target-has-device-addr-8.C: New test.
	* testsuite/libgomp.c++/target-has-device-addr-9.C: New test.

(cherry picked from commit b4fb9f4f9a10d825302cfb5a0ecefa796570d8bc)

diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index 17d8625e00933b2412693b02a59cd457e0674006..5704076ab70c6a854507e298cff3ae3e614d0c9a 100644 (file)
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,11 @@
+2022-07-04  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backport from mainline:
+       2022-05-16  Marcel Vollweiler  <marcel@codesourcery.com>
+
+       * pt.cc (tsubst_omp_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR.
+       * semantics.cc (finish_omp_clauses): Added template decl processing.
+
 2022-06-30  Tobias Burnus  <tobias@codesourcery.com>
 
        * parser.cc (cp_parser_omp_requires): Add missing %<...%> in error.

diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index 3aa7c1e0211ddfe1042fc69618b94b1d86c11e11..f1df95efbc945e0cc520625b6e3e270d8e74468c 100644 (file)
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -17740,6 +17740,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
        case OMP_CLAUSE_USE_DEVICE_PTR:
        case OMP_CLAUSE_USE_DEVICE_ADDR:
        case OMP_CLAUSE_IS_DEVICE_PTR:
+       case OMP_CLAUSE_HAS_DEVICE_ADDR:
        case OMP_CLAUSE_INCLUSIVE:
        case OMP_CLAUSE_EXCLUSIVE:
          OMP_CLAUSE_DECL (nc)
@@ -17885,6 +17886,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
          case OMP_CLAUSE_USE_DEVICE_PTR:
          case OMP_CLAUSE_USE_DEVICE_ADDR:
          case OMP_CLAUSE_IS_DEVICE_PTR:
+         case OMP_CLAUSE_HAS_DEVICE_ADDR:
          case OMP_CLAUSE_INCLUSIVE:
          case OMP_CLAUSE_EXCLUSIVE:
          case OMP_CLAUSE_ALLOCATE:

diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 8a79def9be7ead1edce2814b6adfa8603a8f9bf9..51918d5091c46f815a0eea054429272aa58f228d 100644 (file)
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8759,14 +8759,20 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
              else
                {
                  t = OMP_CLAUSE_DECL (c);
+                 while (TREE_CODE (t) == TREE_LIST)
+                   t = TREE_CHAIN (t);
                  while (TREE_CODE (t) == INDIRECT_REF
                         || TREE_CODE (t) == ARRAY_REF)
                    t = TREE_OPERAND (t, 0);
                }
            }
-         bitmap_set_bit (&is_on_device_head, DECL_UID (t));
          if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
-           cxx_mark_addressable (t);
+           {
+             bitmap_set_bit (&is_on_device_head, DECL_UID (t));
+             if (!processing_template_decl
+                 && !cxx_mark_addressable (t))
+               remove = true;
+           }
          goto check_dup_generic_t;
 
        case OMP_CLAUSE_USE_DEVICE_ADDR:

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 43b749417c83a3f3fc8686267cf1e4456a92f65a..d10eb5174df1ac286e1ad881185a7b1ac7eb9bd8 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,12 @@
+2022-07-04  Tobias Burnus  <tobias@codesourcery.com>
+
+       Backport from mainline:
+       2022-05-16  Marcel Vollweiler  <marcel@codesourcery.com>
+
+       * testsuite/libgomp.c++/target-has-device-addr-7.C: New test.
+       * testsuite/libgomp.c++/target-has-device-addr-8.C: New test.
+       * testsuite/libgomp.c++/target-has-device-addr-9.C: New test.
+
 2022-07-04  Tobias Burnus  <tobias@codesourcery.com>
 
        Backport from mainline:

diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-7.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-7.C
new file mode 100644 (file)
index 0000000..2c4571b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-7.C
@@ -0,0 +1,36 @@
+/* Testing 'has_device_addr' clause on the target construct with template.  */
+
+template <typename T>
+void
+foo (T x)
+{
+  x = 24;
+  #pragma omp target data map(x) use_device_addr(x)
+    #pragma omp target has_device_addr(x)
+      x = 42;
+
+  if (x != 42)
+    __builtin_abort ();
+}
+
+template <typename T>
+void
+bar (T (&x)[])
+{
+  x[0] = 24;
+  #pragma omp target data map(x[:2]) use_device_addr(x)
+    #pragma omp target has_device_addr(x[:2])
+      x[0] = 42;
+
+  if (x[0] != 42)
+    __builtin_abort ();
+}
+
+int
+main ()
+{
+  int a[] = { 24, 42};
+  foo <int> (42);
+  bar <int> (a);
+  return 0;
+}

diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-8.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-8.C
new file mode 100644 (file)
index 0000000..2adfd30
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-8.C
@@ -0,0 +1,47 @@
+/* Testing 'has_device_addr' clause on the target construct with template.  */
+
+#include <omp.h>
+
+template <typename T>
+void
+foo (T (&x)[])
+{
+  #pragma omp target has_device_addr(x)
+    for (int i = 0; i < 15; i++)
+      x[i] = 2 * i;
+
+  #pragma omp target has_device_addr(x[15:15])
+    for (int i = 15; i < 30; i++)
+      x[i] = 3 * i;
+}
+
+int
+main ()
+{
+  int *dp = (int*)omp_target_alloc (30*sizeof(int), 0);
+
+  #pragma omp target is_device_ptr(dp)
+    for (int i = 0; i < 30; i++)
+      dp[i] = i;
+
+  int (&x)[30] = *static_cast<int(*)[30]>(static_cast<void*>(dp));
+
+  foo <int> (x);
+
+  int y[30];
+  for (int i = 0; i < 30; ++i)
+    y[i] = 0;
+  int h = omp_get_initial_device ();
+  int t = omp_get_default_device ();
+  omp_target_memcpy (&y, dp, 30 * sizeof(int), 0, 0, h, t);
+  for (int i = 0; i < 15; ++i)
+    if (y[i] != 2 * i)
+      __builtin_abort ();
+  for (int i = 15; i < 30; ++i)
+    if (y[i] != 3 * i)
+      __builtin_abort ();
+
+  omp_target_free (dp, 0);
+
+  return 0;
+}

diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-9.C b/libgomp/testsuite/libgomp.c++/target-has-device-addr-9.C
new file mode 100644 (file)
index 0000000..0c34cab
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-has-device-addr-9.C
@@ -0,0 +1,30 @@
+/* Testing 'has_device_addr' clause on the target construct with template.  */
+
+#include <omp.h>
+
+template <typename T>
+void
+foo (T (&x))
+{
+  #pragma omp target has_device_addr(x)
+    x = 24;
+}
+
+int
+main ()
+{
+  int *dp = (int*)omp_target_alloc (sizeof(int), 0);
+  int &x = *dp;
+
+  foo <int> (x);
+
+  int y = 42;
+  int h = omp_get_initial_device ();
+  int t = omp_get_default_device ();
+  omp_target_memcpy (&y, dp, sizeof(int), 0, 0, h, t);
+  if (y != 24)
+    __builtin_abort ();
+
+  omp_target_free (dp, 0);
+  return 0;
+}