mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-03-22 00:01:10 +08:00
Code Issues Packages Projects Releases Wiki Activity

openmp: Fix handling of non-addressable shared scalars in parallel nested inside of target [PR93515]

Browse Source 
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  <jakub@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.
This commit is contained in:
Jakub Jelinek 2020-02-06 09:19:08 +01:00
parent cf785618ec
commit cb3f06480a
4 changed files with 73 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
gcc/ChangeLog
View File

@ -1,5 +1,11 @@
2020-02-06 Jakub Jelinek <jakub@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.
PR libgomp/93515
* gimplify.c (gimplify_scan_omp_clauses) <do_notice>: If adding
shared clause, call omp_notice_variable on outer context if any.

33
gcc/omp-low.c
View File

@ -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 <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:

5
libgomp/ChangeLog
View File

@ -1,3 +1,8 @@
2020-02-06 Jakub Jelinek <jakub@redhat.com>
PR libgomp/93515
* testsuite/libgomp.c-c++-common/pr93515.c: New test.
2020-02-05 Tobias Burnus <tobias@codesourcery.com>
* testsuite/lib/libgomp.exp

36
libgomp/testsuite/libgomp.c-c++-common/pr93515.c Normal file
View File

@ -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;
}