Hi!
As the following testcase shows, we need to consider even target to be a
construct
that forces not to use copy in/out for shared on parallel inside of the target.
E.g. for parallel nested inside another parallel or host teams, we already avoid
copy in/out and we need to treat target the same.
Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk.
2020-02-06 Jakub Jelinek <ja...@redhat.com>
PR libgomp/93515
* omp-low.c (use_pointer_for_field): For nested constructs, also
look for map clauses on target construct.
(scan_omp_1_stmt) <case GIMPLE_OMP_TARGET>: Bump temporarily
taskreg_nesting_level.
* testsuite/libgomp.c-c++-common/pr93515.c: New test.
--- gcc/omp-low.c.jj 2020-01-12 11:54:36.688409260 +0100
+++ gcc/omp-low.c 2020-01-31 15:00:46.852168424 +0100
@@ -477,18 +477,30 @@ use_pointer_for_field (tree decl, omp_co
omp_context *up;
for (up = shared_ctx->outer; up; up = up->outer)
- if (is_taskreg_ctx (up) && maybe_lookup_decl (decl, up))
+ if ((is_taskreg_ctx (up)
+ || (gimple_code (up->stmt) == GIMPLE_OMP_TARGET
+ && is_gimple_omp_offloaded (up->stmt)))
+ && maybe_lookup_decl (decl, up))
break;
if (up)
{
tree c;
- for (c = gimple_omp_taskreg_clauses (up->stmt);
- c; c = OMP_CLAUSE_CHAIN (c))
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED
- && OMP_CLAUSE_DECL (c) == decl)
- break;
+ if (gimple_code (up->stmt) == GIMPLE_OMP_TARGET)
+ {
+ for (c = gimple_omp_target_clauses (up->stmt);
+ c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_DECL (c) == decl)
+ break;
+ }
+ else
+ for (c = gimple_omp_taskreg_clauses (up->stmt);
+ c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED
+ && OMP_CLAUSE_DECL (c) == decl)
+ break;
if (c)
goto maybe_mark_addressable_and_ret;
@@ -3781,7 +3793,14 @@ scan_omp_1_stmt (gimple_stmt_iterator *g
break;
case GIMPLE_OMP_TARGET:
- scan_omp_target (as_a <gomp_target *> (stmt), ctx);
+ if (is_gimple_omp_offloaded (stmt))
+ {
+ taskreg_nesting_level++;
+ scan_omp_target (as_a <gomp_target *> (stmt), ctx);
+ taskreg_nesting_level--;
+ }
+ else
+ scan_omp_target (as_a <gomp_target *> (stmt), ctx);
break;
case GIMPLE_OMP_TEAMS:
--- libgomp/testsuite/libgomp.c-c++-common/pr93515.c.jj 2020-01-31
14:53:01.163112148 +0100
+++ libgomp/testsuite/libgomp.c-c++-common/pr93515.c 2020-01-31
14:52:38.627448474 +0100
@@ -0,0 +1,36 @@
+/* PR libgomp/93515 */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+ int i;
+ int a = 42;
+#pragma omp target teams distribute parallel for defaultmap(tofrom: scalar)
+ for (i = 0; i < 64; ++i)
+ if (omp_get_team_num () == 0)
+ if (omp_get_thread_num () == 0)
+ a = 142;
+ if (a != 142)
+ __builtin_abort ();
+ a = 42;
+#pragma omp target parallel for defaultmap(tofrom: scalar)
+ for (i = 0; i < 64; ++i)
+ if (omp_get_thread_num () == 0)
+ a = 143;
+ if (a != 143)
+ __builtin_abort ();
+ a = 42;
+#pragma omp target firstprivate(a)
+ {
+ #pragma omp parallel for
+ for (i = 0; i < 64; ++i)
+ if (omp_get_thread_num () == 0)
+ a = 144;
+ if (a != 144)
+ abort ();
+ }
+ return 0;
+}
Jakub
