* [PATCH] OpenMP, C++: Add template support for the has_device_addr clause.
@ 2022-02-23 16:01 Marcel Vollweiler
2022-05-05 14:06 ` Jakub Jelinek
0 siblings, 1 reply; 4+ messages in thread
From: Marcel Vollweiler @ 2022-02-23 16:01 UTC (permalink / raw)
To: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 563 bytes --]
Hi,
The patch for adding the has_device_addr clause on the target construct was
recently committed (bbb7f8604e1dfc08f44354cfd93d2287f2fdd489).
Additionally, this patch adds support for list items in the has_device_addr
clause which type is given by C++ template parameters.
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
[-- Attachment #2: has-device-addr-template-patch.diff --]
[-- Type: text/plain, Size: 8107 bytes --]
OpenMP, C++: Add template support for the has_device_addr clause.
gcc/cp/ChangeLog:
* pt.cc (tsubst_omp_clauses): Add OMP_CLAUSE_HAS_DEVICE_ADDR.
* semantics.cc (finish_omp_clauses): Handle PARM_DECL and
NON_LVALUE_EXPR.
gcc/ChangeLog:
* gimplify.cc (gimplify_scan_omp_clauses): Handle NON_LVALUE_EXPR.
(gimplify_adjust_omp_clauses): Likewise.
* omp-low.cc (scan_sharing_clauses): Likewise.
(lower_omp_target): Likewise.
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 6dda660..86446d7 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -17652,6 +17652,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)
@@ -17797,6 +17798,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 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);
}
}
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);
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;
+}
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH] OpenMP, C++: Add template support for the has_device_addr clause.
2022-02-23 16:01 [PATCH] OpenMP, C++: Add template support for the has_device_addr clause Marcel Vollweiler
@ 2022-05-05 14:06 ` Jakub Jelinek
2022-05-10 15:41 ` Marcel Vollweiler
0 siblings, 1 reply; 4+ messages in thread
From: Jakub Jelinek @ 2022-05-05 14:06 UTC (permalink / raw)
To: Marcel Vollweiler; +Cc: gcc-patches
On Wed, Feb 23, 2022 at 05:01:45PM +0100, Marcel Vollweiler wrote:
> gcc/cp/ChangeLog:
>
> * pt.cc (tsubst_omp_clauses): Add OMP_CLAUSE_HAS_DEVICE_ADDR.
> * semantics.cc (finish_omp_clauses): Handle PARM_DECL and
> NON_LVALUE_EXPR.
>
> gcc/ChangeLog:
>
> * gimplify.cc (gimplify_scan_omp_clauses): Handle NON_LVALUE_EXPR.
> (gimplify_adjust_omp_clauses): Likewise.
> * omp-low.cc (scan_sharing_clauses): Likewise.
> (lower_omp_target): Likewise.
>
> 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.
>
> --- a/gcc/cp/pt.cc
> +++ b/gcc/cp/pt.cc
> @@ -17652,6 +17652,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)
> @@ -17797,6 +17798,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:
This part is ok.
> 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.
> 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...
Jakub
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH] OpenMP, C++: Add template support for the has_device_addr clause.
2022-05-05 14:06 ` Jakub Jelinek
@ 2022-05-10 15:41 ` Marcel Vollweiler
2022-05-10 15:44 ` Jakub Jelinek
0 siblings, 1 reply; 4+ messages in thread
From: Marcel Vollweiler @ 2022-05-10 15:41 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 8208 bytes --]
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
[-- Attachment #2: has-device-addr-template-patch.diff --]
[-- Type: text/plain, Size: 5006 bytes --]
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;
+}
^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH] OpenMP, C++: Add template support for the has_device_addr clause.
2022-05-10 15:41 ` Marcel Vollweiler
@ 2022-05-10 15:44 ` Jakub Jelinek
0 siblings, 0 replies; 4+ messages in thread
From: Jakub Jelinek @ 2022-05-10 15:44 UTC (permalink / raw)
To: Marcel Vollweiler; +Cc: gcc-patches
On Tue, May 10, 2022 at 05:41:38PM +0200, Marcel Vollweiler wrote:
> 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.
Ok, thanks.
> 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.
Jakub
^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2022-05-13 15:21 UTC | newest]
Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-02-23 16:01 [PATCH] OpenMP, C++: Add template support for the has_device_addr clause Marcel Vollweiler
2022-05-05 14:06 ` Jakub Jelinek
2022-05-10 15:41 ` Marcel Vollweiler
2022-05-10 15:44 ` Jakub Jelinek
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).