From cb3f06480a17f98579704b9927632627a3814c5c Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Thu, 6 Feb 2020 09:19:08 +0100 Subject: [PATCH] openmp: Fix handling of non-addressable shared scalars in parallel nested inside of target [PR93515] 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. 2020-02-06 Jakub Jelinek 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) : Bump temporarily taskreg_nesting_level. * testsuite/libgomp.c-c++-common/pr93515.c: New test. --- gcc/ChangeLog | 6 ++++ gcc/omp-low.c | 33 +++++++++++++++++----- libgomp/ChangeLog | 5 ++++ libgomp/testsuite/libgomp.c-c++-common/pr93515.c | 36 ++++++++++++++++++++++++ 4 files changed, 73 insertions(+), 7 deletions(-) create mode 100644 libgomp/testsuite/libgomp.c-c++-common/pr93515.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index df6bbe6..b5b465a 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,6 +1,12 @@ 2020-02-06 Jakub Jelinek 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) : Bump temporarily + taskreg_nesting_level. + + PR libgomp/93515 * gimplify.c (gimplify_scan_omp_clauses) : If adding shared clause, call omp_notice_variable on outer context if any. diff --git a/gcc/omp-low.c b/gcc/omp-low.c index eb3fe96..67565d6 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -477,18 +477,30 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx) 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 *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: - scan_omp_target (as_a (stmt), ctx); + if (is_gimple_omp_offloaded (stmt)) + { + taskreg_nesting_level++; + scan_omp_target (as_a (stmt), ctx); + taskreg_nesting_level--; + } + else + scan_omp_target (as_a (stmt), ctx); break; case GIMPLE_OMP_TEAMS: diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 56e9da7..ba14005 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2020-02-06 Jakub Jelinek + + PR libgomp/93515 + * testsuite/libgomp.c-c++-common/pr93515.c: New test. + 2020-02-05 Tobias Burnus * testsuite/lib/libgomp.exp diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr93515.c b/libgomp/testsuite/libgomp.c-c++-common/pr93515.c new file mode 100644 index 0000000..8a69088 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/pr93515.c @@ -0,0 +1,36 @@ +/* PR libgomp/93515 */ + +#include +#include + +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; +} -- 2.7.4