* [patch] OpenMP: Handle same-directive mapped vars with pointer predefined firstprivate [PR110639]
@ 2023-12-08 14:28 Tobias Burnus
2023-12-08 14:43 ` Jakub Jelinek
0 siblings, 1 reply; 2+ messages in thread
From: Tobias Burnus @ 2023-12-08 14:28 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 3082 bytes --]
This patch fixes the issue:
int a[100];
p = &a[0];
#pragma omp target map(a)
p[0] = p[99] = 3;
where 'p' is predetermined firstprivate, i.e. it is firstprivatized
but its address gets updated to the device address of 'a' as there is
associated storage for the value of 'p', i.e. its pointee.
[This is a C/C++-only feature that cannot be replicated by using a single clause.
('target data map(a) use_device_ptr(p)' + 'target is_device_ptr(p)' would do
so in two steps. - or 'p2 = omp_get_mapped_ptr(p, devnum)' + 'target is_device_ptr(p2)'.)]
Before this only worked when that storage was mapped before and not on the same
directive.
The gimplify_scan_omp_clauses change was done when I saw some runtime fails; I think
those were due to a bug in libgomp (now fixed) and not due to having two pointer
privatisations in a now different order. Still, they at least prevent mapping
'this' multiple times when 'this' is not 'this' but __closure->this which is at least
a missed optimization. And also for libgomp.c++/pr108286.C which has a normal
'this' and map(tofrom:*this [len: 16]).
Build and tested without offloading and with nvptx offloading.
Comments, remarks, suggestions?
* * *
(I wonder whether our current approach of removing explicit MAP if its
DECL is unsued is the right one if there is any GOVD_MAP_0LEN_ARRAY around
- or even any OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION.)
(See new libgomp.c-c++-common/target-implicit-map-6.c; BTW, I tried:
before '(void) a;' but that only worked with C and not with C++.)
* * *
The other issue in the PR (still to be done) is for code like:
int a[100];
p = &a[0];
#pragma omp target map(a[20:20]) // Map only a[20] to a[40], but p points to &a[0]
p[20] = p[30] = 3;
where 'p' points to the base address of 'a' but p[0] == a[0] it not actually
mapped. As we currently do not keep track of base pointer, this won't work.
I have not (yet) explored how to best implement this.
* * *
OpenMP Spec:
The first feature is not new, but I have not checked the wording in 4.5 or 5.0;
it might be that older versions only required it to work for storage mapped before
the current taget directive. But at least TR12 is very explicit in permitting it
and the (nonpublic) issue which lead to the 5.1 change also uses this. (See PR.)
(The second feature is definitely new in OpenMP 5.1.)
TR12 states in "14.8 target Construct" [379:8-10]:
"[C/C++] If a list item in a map clause has a base pointer that is predetermined firstprivate
(see Section 6.1.1) and on entry to the target region the list item is mapped, the firstprivate
pointer is updated via corresponding base pointer initialization."
(For OpenMP 5.1, read its Section 2.21.7.2.)
Tobias
-----------------
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: map-firstprivate.diff --]
[-- Type: text/x-patch, Size: 19802 bytes --]
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(-)
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 342e43a7f25..2234fd6b7e1 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -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);
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
index 5ce8ceadb19..b4f1593af02 100644
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -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" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C
index cc08e7e8693..eecab5a25e8 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-2.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C
@@ -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" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index 9ade3cc0b2b..845f89f6997 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -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" } } */
diff --git a/libgomp/target.c b/libgomp/target.c
index f30c20255d3..c1f26e6860b 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -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'
diff --git a/libgomp/testsuite/libgomp.c++/pr108286.C b/libgomp/testsuite/libgomp.c++/pr108286.C
index ee88c2f9fd0..3ae2f5dd16e 100644
--- a/libgomp/testsuite/libgomp.c++/pr108286.C
+++ b/libgomp/testsuite/libgomp.c++/pr108286.C
@@ -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" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-6.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-6.c
new file mode 100644
index 00000000000..6a2d3bfd0e3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-6.c
@@ -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;
+}
^ permalink raw reply [flat|nested] 2+ messages in thread
* Re: [patch] OpenMP: Handle same-directive mapped vars with pointer predefined firstprivate [PR110639]
2023-12-08 14:28 [patch] OpenMP: Handle same-directive mapped vars with pointer predefined firstprivate [PR110639] Tobias Burnus
@ 2023-12-08 14:43 ` Jakub Jelinek
0 siblings, 0 replies; 2+ messages in thread
From: Jakub Jelinek @ 2023-12-08 14:43 UTC (permalink / raw)
To: Tobias Burnus; +Cc: gcc-patches
On Fri, Dec 08, 2023 at 03:28:59PM +0100, Tobias Burnus wrote:
> This patch fixes the issue:
>
> int a[100];
> p = &a[0];
>
> #pragma omp target map(a)
> p[0] = p[99] = 3;
>
> where 'p' is predetermined firstprivate, i.e. it is firstprivatized
> but its address gets updated to the device address of 'a' as there is
> associated storage for the value of 'p', i.e. its pointee.
I think the above is invalid even in TR12.
> OpenMP Spec:
>
> The first feature is not new, but I have not checked the wording in 4.5 or 5.0;
> it might be that older versions only required it to work for storage mapped before
> the current taget directive. But at least TR12 is very explicit in permitting it
> and the (nonpublic) issue which lead to the 5.1 change also uses this. (See PR.)
> (The second feature is definitely new in OpenMP 5.1.)
>
> TR12 states in "14.8 target Construct" [379:8-10]:
>
> "[C/C++] If a list item in a map clause has a base pointer that is predetermined firstprivate
> (see Section 6.1.1) and on entry to the target region the list item is mapped, the firstprivate
> pointer is updated via corresponding base pointer initialization."
The list item (a) in the above case doesn't have a base pointer, but base
array. See the glossary. So, the rule would be about something like
int *p = ...;
#pragma omp target map (p[20:100]) or similar, not about an array and an
unrelated pointer.
Jakub
^ permalink raw reply [flat|nested] 2+ messages in thread
end of thread, other threads:[~2023-12-08 14:43 UTC | newest]
Thread overview: 2+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-12-08 14:28 [patch] OpenMP: Handle same-directive mapped vars with pointer predefined firstprivate [PR110639] Tobias Burnus
2023-12-08 14:43 ` 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).