Hi Jakub,
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 0cb17a6..452ecfd 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8534,11 +8534,14 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
{
if (handle_omp_array_sections (c, ort))
remove = true;
+ else if (TREE_CODE (TREE_CHAIN (t)) == PARM_DECL)
+ t = TREE_CHAIN (t);
else
{
t = OMP_CLAUSE_DECL (c);
while (TREE_CODE (t) == INDIRECT_REF
- || TREE_CODE (t) == ARRAY_REF)
+ || TREE_CODE (t) == ARRAY_REF
+ || TREE_CODE (t) == NON_LVALUE_EXPR)
t = TREE_OPERAND (t, 0);
}
}
This is wrong.
When processing_template_decl, handle_omp_array_sections often punts, keeps
things as is because if something is dependent, we can't do much about it.
The else if (TREE_CODE (TREE_CHAIN (t)) == PARM_DECL) is obviously wrong,
there is really nothing specific about PARM_DECLs (just that you used
exactly that in the testcase), nor about array section with exactly one
dimension. What is done elsewhere is look through all TREE_LISTs to find
the base expression, and if that expression is a VAR_DECL/PARM_DECL, nice,
we can do further processing, if processing_template_decl and it is
something different, just defer and otherwise error out.
So I think you want:
--- gcc/cp/semantics.cc.jj 2022-05-05 11:56:16.160443828 +0200
+++ gcc/cp/semantics.cc 2022-05-05 15:52:39.651211448 +0200
@@ -8553,14 +8553,23 @@ finish_omp_clauses (tree clauses, enum c
else
{
t = OMP_CLAUSE_DECL (c);
+ if (TREE_CODE (t) == TREE_LIST)
+ {
+ 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:
instead, as I said look through the TREE_LISTs, then only use DECL_UID
on actual VAR_DECLs/PARM_DECLs not random other expressions and
never call cxx_mark_addressable when processing_template_decl (and remove
clause if cxx_mark_addressable fails).
Note, check_dup_generic_t will do among other things:
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL
&& (!field_ok || TREE_CODE (t) != FIELD_DECL))
{
if (processing_template_decl && TREE_CODE (t) != OVERLOAD)
break;
... error ...
}
so with processing_template_decl it will just defer it for later,
but otherwise if t is something invalid it will diagnose it.
But one really shouldn't rely on t being VAR_DECL/PARM_DECL before
that checking is done...
With your pt.cc change and my semantics.cc change, all your new testcases
look fine.
Thank you very much for your detailed explanation. That helped me a lot for my
understanding!
I adjusted the code accordingly.
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index f570daa..b1bb5be 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10285,7 +10285,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
case OMP_CLAUSE_HAS_DEVICE_ADDR:
decl = OMP_CLAUSE_DECL (c);
while (TREE_CODE (decl) == INDIRECT_REF
- || TREE_CODE (decl) == ARRAY_REF)
+ || TREE_CODE (decl) == ARRAY_REF
+ || TREE_CODE (decl) == NON_LVALUE_EXPR)
decl = TREE_OPERAND (decl, 0);
flags = GOVD_EXPLICIT;
goto do_add_decl;
@@ -11443,7 +11444,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
case OMP_CLAUSE_HAS_DEVICE_ADDR:
decl = OMP_CLAUSE_DECL (c);
while (TREE_CODE (decl) == INDIRECT_REF
- || TREE_CODE (decl) == ARRAY_REF)
+ || TREE_CODE (decl) == ARRAY_REF
+ || TREE_CODE (decl) == NON_LVALUE_EXPR)
decl = TREE_OPERAND (decl, 0);
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
remove = n == NULL || !(n->value & GOVD_SEEN);
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 77176ef..30cc9b6 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1384,7 +1384,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
- if (TREE_CODE (decl) == INDIRECT_REF)
+ if (TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == NON_LVALUE_EXPR)
decl = TREE_OPERAND (decl, 0);
install_var_field (decl, true, 3, ctx);
}
@@ -1747,7 +1748,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
while (TREE_CODE (decl) == INDIRECT_REF
- || TREE_CODE (decl) == ARRAY_REF)
+ || TREE_CODE (decl) == ARRAY_REF
+ || TREE_CODE (decl) == NON_LVALUE_EXPR)
decl = TREE_OPERAND (decl, 0);
}
@@ -12847,7 +12849,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
while (TREE_CODE (var) == INDIRECT_REF
- || TREE_CODE (var) == ARRAY_REF)
+ || TREE_CODE (var) == ARRAY_REF
+ || TREE_CODE (var) == NON_LVALUE_EXPR)
var = TREE_OPERAND (var, 0);
}
map_cnt++;
@@ -13337,7 +13340,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
while (TREE_CODE (ovar) == INDIRECT_REF
- || TREE_CODE (ovar) == ARRAY_REF)
+ || TREE_CODE (ovar) == ARRAY_REF
+ || TREE_CODE (ovar) == NON_LVALUE_EXPR)
ovar = TREE_OPERAND (ovar, 0);
}
var = lookup_decl_in_outer_ctx (ovar, ctx);
@@ -13607,7 +13611,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
while (TREE_CODE (var) == INDIRECT_REF
- || TREE_CODE (var) == ARRAY_REF)
+ || TREE_CODE (var) == ARRAY_REF
+ || TREE_CODE (var) == NON_LVALUE_EXPR)
var = TREE_OPERAND (var, 0);
}
x = build_receiver_ref (var, false, ctx);
So all the gimplify.cc and omp-low.cc changes look suspicious.
NON_LVALUE_EXPR shouldn't show up there...
I removed all the NON_LVALUE_EXPR changes again.
The new version of the patch was tested again on x86_64-linux with nvptx and
amdgcn offloading. All with no regressions.
Marcel
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
OpenMP, C++: Add template support for the has_device_addr clause.
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.
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index fe2608c..a07e5a4 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -17722,6 +17722,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)
@@ -17867,6 +17868,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 10478d1..1542aed 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8553,14 +8553,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/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 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
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
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;
+}
