OpenMP: Handle same-directive mapped vars with pointer predefined firstprivate [PR110639]
Predefined 'firstprivate' for pointer variables firstprivatizes the pointer
but if it is associated with a mapped target, its address is updated to the
corresponding target. (If not, the host value remains.)
This commit extends this handling to also update the pointer address for
storaged mapped on the same directive.
The 'gimplify_scan_omp_clauses' change avoids adding an additional
map(alloc:this) (+ptr assignment)
when there is already a
map(tofrom:*this) (+ptr assignment)
This shows up for libgomp.c++/pr108286.C and also when 'this' is
actually '__closure->this' (-> g++.dg/gomp/target-{this-{2,4},lambda-1}.C).
PR middle-end/110639
gcc/ChangeLog:
* gimplify.cc (struct gimplify_adjust_omp_clauses_data): Add
append_list.
(gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Add
GOVD_MAP_0LEN_ARRAY clauses at the end.
(gimplify_scan_omp_clauses): Mark also '*var' as found not only
'var'.
libgomp/ChangeLog:
* target.c (gomp_map_vars_internal): Handle also variables
mapped in the same directive for GOVD_MAP_0LEN_ARRAY.
* testsuite/libgomp.c++/pr108286.C: Add gimple tree-scan test.
* testsuite/libgomp.c-c++-common/target-implicit-map-6.c: New test.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/target-this-2.C: Remove 'this' pointer mapping alreay
mapped via __closure->this.
* g++.dg/gomp/target-this-4.C: Likewise.
* g++.dg/gomp/target-lambda-1.C: Likewise. Move 'iptr' pointer
mapping to the end in scan-tree-dump.
gcc/gimplify.cc | 45 ++++-
gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 4 +-
gcc/testsuite/g++.dg/gomp/target-this-2.C | 4 +-
gcc/testsuite/g++.dg/gomp/target-this-4.C | 6 +-
libgomp/target.c | 11 +-
libgomp/testsuite/libgomp.c++/pr108286.C | 4 +
.../libgomp.c-c++-common/target-implicit-map-6.c | 212 +++++++++++++++++++++
7 files changed, 276 insertions(+), 10 deletions(-)
@@ -11586,6 +11586,23 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
else if (!DECL_P (decl))
{
tree d = decl, *pd;
+ pd = &OMP_CLAUSE_DECL (c);
+ if (TREE_CODE (decl) == INDIRECT_REF)
+ {
+ tree d2 = TREE_OPERAND (decl, 0);
+ STRIP_NOPS (d2);
+ if (DECL_P (d2))
+ {
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+ fb_lvalue) == GS_ERROR)
+ {
+ remove = true;
+ break;
+ }
+ decl = d2;
+ goto handle_map_decl;
+ }
+ }
if (TREE_CODE (d) == ARRAY_REF)
{
while (TREE_CODE (d) == ARRAY_REF)
@@ -11594,7 +11611,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
&& TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE)
decl = d;
}
- pd = &OMP_CLAUSE_DECL (c);
if (d == decl
&& TREE_CODE (decl) == INDIRECT_REF
&& TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
@@ -11774,6 +11790,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
break;
}
+ handle_map_decl:
flags = GOVD_MAP | GOVD_EXPLICIT;
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
@@ -11806,7 +11823,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
}
- goto do_add;
+ goto do_add_decl;
case OMP_CLAUSE_AFFINITY:
gimplify_omp_affinity (list_p, pre_p);
@@ -12571,6 +12588,7 @@ omp_find_stores_stmt (gimple_stmt_iterator *gsi_p,
struct gimplify_adjust_omp_clauses_data
{
tree *list_p;
+ tree append_list;
gimple_seq *pre_p;
};
@@ -12691,6 +12709,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
&& omp_shared_to_firstprivate_optimizable_decl_p (decl))
omp_mark_stores (gimplify_omp_ctxp->outer_context, decl);
+ bool len0_append_list_used = false;
tree chain = *list_p;
clause = build_omp_clause (input_location, code);
OMP_CLAUSE_DECL (clause) = decl;
@@ -12707,6 +12726,11 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (clause) = 1;
else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0)
{
+ /* For GOVD_MAP_0LEN_ARRAY, add the clauses to append_list such
+ that those come after any data mapping. */
+ len0_append_list_used = true;
+ struct gimplify_adjust_omp_clauses_data *adjdata
+ = (struct gimplify_adjust_omp_clauses_data *) data;
tree nc = build_omp_clause (input_location, OMP_CLAUSE_MAP);
OMP_CLAUSE_DECL (nc) = decl;
if (TREE_CODE (TREE_TYPE (decl)) == REFERENCE_TYPE
@@ -12721,8 +12745,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_ALLOC);
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (clause) = 1;
OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_POINTER);
- OMP_CLAUSE_CHAIN (nc) = chain;
+ OMP_CLAUSE_CHAIN (nc) = adjdata->append_list;
OMP_CLAUSE_CHAIN (clause) = nc;
+ adjdata->append_list = clause;
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
gimplify_omp_ctxp = ctx->outer_context;
gimplify_expr (&TREE_OPERAND (OMP_CLAUSE_DECL (clause), 0),
@@ -12833,7 +12858,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
(ctx->region_type & ORT_ACC) != 0);
gimplify_omp_ctxp = ctx;
}
- *list_p = clause;
+ if (!len0_append_list_used)
+ *list_p = clause;
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
gimplify_omp_ctxp = ctx->outer_context;
/* Don't call omp_finish_clause on implicitly added OMP_CLAUSE_PRIVATE
@@ -12842,7 +12868,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
if (code != OMP_CLAUSE_PRIVATE || ctx->region_type != ORT_SIMD)
lang_hooks.decls.omp_finish_clause (clause, pre_p,
(ctx->region_type & ORT_ACC) != 0);
- if (gimplify_omp_ctxp)
+ if (gimplify_omp_ctxp && !len0_append_list_used)
for (; clause != chain; clause = OMP_CLAUSE_CHAIN (clause))
if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
&& DECL_P (OMP_CLAUSE_SIZE (clause)))
@@ -13445,6 +13471,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
/* Add in any implicit data sharing. */
struct gimplify_adjust_omp_clauses_data data;
+ data.append_list = NULL;
if ((gimplify_omp_ctxp->region_type & ORT_ACC) == 0)
{
/* OpenMP. Implicit clauses are added at the start of the clause list,
@@ -13472,6 +13499,14 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
"iterator");
break;
}
+ if (data.append_list != NULL_TREE && *data.list_p != NULL_TREE)
+ {
+ for (c = *data.list_p; c && OMP_CLAUSE_CHAIN (c); c = OMP_CLAUSE_CHAIN (c))
+ ;
+ OMP_CLAUSE_CHAIN (c) = data.append_list;
+ }
+ else if (data.append_list != NULL_TREE)
+ *data.list_p = data.append_list;
gimplify_omp_ctxp = ctx->outer_context;
delete_omp_context (ctx);
@@ -87,7 +87,9 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
+/* Note that 'this' = '__closure->__this' such that no pointer-assign for 'this' should appear. */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\)[\r\n]} "gimple" } } */
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */
@@ -46,4 +46,6 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {map\(alloc:MEM\[\(char \*\)_[0-9]+\] \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) firstprivate\(m\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:v \[len: [0-9]+\]\)} "gimple" } } */
+/* Note that 'this' = '__closure->__this' such that no pointer-assign for 'this' should appear. */
+
+/* { dg-final { scan-tree-dump {firstprivate\(n\) firstprivate\(m\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:v \[len: [0-9]+\]\)[\r\n]} "gimple" } } */
@@ -102,6 +102,8 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
+/* Note that 'this' = '__closure->__this' such that no pointer-assign for 'this' should appear. */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)[\r\n]} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(n\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\)[\r\n]} "gimple" } } */
@@ -1149,7 +1149,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
splay_tree_key n;
if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
{
- n = gomp_map_0len_lookup (mem_map, &cur_node);
+ /* Defer lookup when mapped item found. */
+ n = not_found_cnt ? NULL : gomp_map_0len_lookup (mem_map, &cur_node);
if (!n)
{
tgt->list[i].key = NULL;
@@ -1417,7 +1418,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
continue;
case GOMP_MAP_FIRSTPRIVATE_INT:
+ continue;
case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizes[i];
+ n = gomp_map_0len_lookup (mem_map, &cur_node);
+ if (n)
+ gomp_map_vars_existing (devicep, aq, n, &cur_node,
+ &tgt->list[i], kind & typemask, false,
+ implicit, NULL, refcount_set);
continue;
case GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT:
/* The OpenACC 'host_data' construct only allows 'use_device'
@@ -1,5 +1,6 @@
// PR c++/108286
// { dg-do run }
+// { dg-additional-options "-fdump-tree-gimple" }
struct S {
int
@@ -27,3 +28,6 @@ main ()
if (s.foo () != 42)
__builtin_abort ();
}
+
+/* Ensure that 'this' is mapped but only once and not additionally via 'this[:0]'. */
+/* { dg-final { scan-tree-dump "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) map\\(tofrom:\\*this \\\[len: \[0-9\]+\\\]\\) map\\(firstprivate:this \\\[pointer assign, bias: 0\\\]\\) nowait map\\(tofrom:res \\\[len: \[0-9\]+\\\]\\) map\\(tofrom:\\*_\[0-9\]+ \\\[len: _\[0-9\]+\\\]\\) map\\(attach:this->ptr \\\[bias: 0\\\]\\)\[\r\n\]" "gimple" } } */
new file mode 100644
@@ -0,0 +1,212 @@
+/* Prefined firstprivate privatizes the pointer
+ and then updates the value to point to the corresponding
+ device variable, if existing.
+
+ See PR middle-end/110639
+ and TR12 in "14.8 target Construct" [379:8-10]
+ or OpenMP 5.1 in "2.21.7.2 Pointer Initialization for Device Data Environments". */
+
+#include <stdlib.h>
+#include <omp.h>
+
+int my_false = 0;
+
+int
+f (int x, int y)
+{
+ return x + y;
+}
+
+void
+no_other_clause ()
+{
+ int data = 3;
+ int *p = &data;
+ #pragma omp target enter data map(data)
+ #pragma omp target
+ *p = 5;
+ #pragma omp target exit data map(data)
+ if (*p != 5)
+ abort ();
+}
+
+void
+test1 (int devnum)
+{
+ int start = 0, n = 100;
+ int a[100];
+ int *p = &a[0];
+
+ for (int i = start; i < start+n; i++)
+ a[i] = 10*i;
+
+ #pragma omp target map(a) device(device_num : devnum)
+ {
+ if (my_false) /* Ensure that 'map(a)' is not optimized away. */
+ a[8] = 1;
+ for (int i = start; i < start+n; i++)
+ p[i] = f(p[i], i);
+ p = NULL;
+ }
+
+ if (p != &a[0])
+ abort ();
+ for (int i = start; i < start+n; i++)
+ if (a[i] != f(10 *i, i))
+ abort ();
+}
+
+
+
+void
+test2 (int devnum)
+{
+ int start = 0, n = 100;
+ int a[100];
+ int *p = &a[0];
+
+ for (int i = start; i < start+n; i++)
+ a[i] = 10*i;
+
+ #pragma omp target enter data map(a) device(device_num : devnum)
+ #pragma omp target device(device_num : devnum)
+ {
+ (void) a; /* Ensure that 'map(a)' is not optimized away. */
+ for (int i = start; i < start+n; i++)
+ p[i] = f(p[i], i);
+ p = NULL;
+ }
+ #pragma omp target exit data map(a) device(device_num : devnum)
+
+ if (p != &a[0])
+ abort ();
+ for (int i = start; i < start+n; i++)
+ if (a[i] != f(10 *i, i))
+ abort ();
+}
+
+void
+test3 (int devnum)
+{
+ int start = 8, n = 10;
+ int a[100];
+ int *p = &a[start];
+
+ for (int i = start; i < start+n; i++)
+ a[i] = 10*i;
+
+ /* p points to a[start] */
+ #pragma omp target map(a[start:n]) device(device_num : devnum)
+ {
+ if (my_false) /* Ensure that 'map(a)' is not optimized away. */
+ a[8] = 1;
+ for (int i = 0; i < n; i++)
+ p[i] = f(p[i], i + start);
+ p = NULL;
+ }
+
+ if (p != &a[start])
+ abort ();
+ for (int i = start; i < start+n; i++)
+ if (a[i] != f(10 *i, i))
+ abort ();
+}
+
+void
+test4 (int devnum)
+{
+ int start = 8, n = 10;
+ int a[100];
+ int *p = &a[start];
+
+ for (int i = start; i < start+n; i++)
+ a[i] = 10*i;
+
+ /* p points to a[start] */
+ #pragma omp target enter data map(a[start:n]) device(device_num : devnum)
+ #pragma omp target device(device_num : devnum)
+ {
+ for (int i = 0; i < n; i++)
+ p[i] = f(p[i], i + start);
+ p = NULL;
+ }
+ #pragma omp target exit data map(a[start:n]) device(device_num : devnum)
+
+ if (p != &a[start])
+ abort ();
+ for (int i = start; i < start+n; i++)
+ if (a[i] != f(10 *i, i))
+ abort ();
+}
+
+void
+test5 (int devnum)
+{
+ int start = 8, n = 10;
+ int a[100];
+ int *p = &a[start + 5];
+
+ for (int i = start; i < start+n; i++)
+ a[i] = 10*i;
+
+ /* p points to a[start + 5] */
+ #pragma omp target map(a[start:n]) device(device_num : devnum)
+ {
+ if (my_false) /* Ensure that 'map(a)' is not optimized away. */
+ a[8] = 1;
+ for (int i = 0; i < n; i++)
+ p[i - 5] = f(p[i - 5], i + start);
+ p = NULL;
+ }
+
+ if (p != &a[start + 5])
+ abort ();
+ for (int i = start; i < start+n; i++)
+ if (a[i] != f(10 *i, i))
+ abort ();
+}
+
+void
+test6 (int devnum)
+{
+ int start = 8, n = 10;
+ int a[100];
+ int *p = &a[start + 5];
+
+ for (int i = start; i < start+n; i++)
+ a[i] = 10*i;
+
+ /* p points to a[start + 5] */
+ #pragma omp target enter data map(a[start:n]) device(device_num : devnum)
+ #pragma omp target device(device_num : devnum)
+ {
+ for (int i = 0; i < n; i++)
+ p[i - 5] = f(p[i - 5], i + start);
+ p = NULL;
+ }
+ #pragma omp target exit data map(a[start:n]) device(device_num : devnum)
+
+ if (p != &a[start + 5])
+ abort ();
+ for (int i = start; i < start+n; i++)
+ if (a[i] != f(10 *i, i))
+ abort ();
+}
+
+int
+main ()
+{
+ int n = omp_get_num_devices ();
+ no_other_clause ();
+ for (int i = omp_initial_device; i <= n; i++)
+ {
+ /* First with 'a' mapped on target; then 'a' on target enter data. */
+ test1 (i); /* p = &a[0], map(a) */
+ test2 (i);
+ test3 (i); /* p = &a[start], map(a[start:n]) */
+ test4 (i);
+ test5 (i); /* p = &a[start + 5], map(a[start:n]) */
+ test6 (i);
+ }
+ return 0;
+}