public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-12] OpenMP, C++: Add template support for the has_device_addr clause.
@ 2022-07-04 20:11 Tobias Burnus
0 siblings, 0 replies; only message in thread
From: Tobias Burnus @ 2022-07-04 20:11 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:b584b32b1f76a9011d5a2473acc2781c78a0802c
commit b584b32b1f76a9011d5a2473acc2781c78a0802c
Author: Marcel Vollweiler <marcel@codesourcery.com>
Date: Mon Jul 4 21:13:39 2022 +0200
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:
---
gcc/cp/ChangeLog.omp | 8 ++++
gcc/cp/pt.cc | 2 +
gcc/cp/semantics.cc | 10 ++++-
libgomp/ChangeLog.omp | 9 +++++
.../libgomp.c++/target-has-device-addr-7.C | 36 +++++++++++++++++
.../libgomp.c++/target-has-device-addr-8.C | 47 ++++++++++++++++++++++
.../libgomp.c++/target-has-device-addr-9.C | 30 ++++++++++++++
7 files changed, 140 insertions(+), 2 deletions(-)
diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index 17d8625e009..5704076ab70 100644
--- 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 3aa7c1e0211..f1df95efbc9 100644
--- 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 8a79def9be7..51918d5091c 100644
--- 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 43b749417c8..d10eb5174df 100644
--- 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
index 00000000000..2c4571be455
--- /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
index 00000000000..2adfd30296f
--- /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
index 00000000000..0c34cab56d2
--- /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;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2022-07-04 20:11 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-07-04 20:11 [gcc/devel/omp/gcc-12] OpenMP, C++: Add template support for the has_device_addr clause Tobias Burnus
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).