OpenE2K/gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[OpenACC] Elaborate/simplify 'exit data' 'finalize' handling

Browse Source 
No functional changes.

	gcc/
	* gimplify.c (gimplify_omp_target_update): Elaborate 'exit data'
	'finalize' handling.
	gcc/testsuite/
	* c-c++-common/goacc/finalize-1.c: Extend.
	* gfortran.dg/goacc/finalize-1.f: Likewise.
	libgomp/
	* oacc-mem.c (GOACC_enter_exit_data): Simplify 'exit data'
	'finalize' handling.

From-SVN: r279531
This commit is contained in:
Thomas Schwinge 2019-12-18 18:00:51 +01:00 committed by Thomas Schwinge
parent ba40277f6a
commit 32128577ae
7 changed files with 49 additions and 22 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
gcc/ChangeLog
View File

@ -1,3 +1,8 @@
2019-12-18 Thomas Schwinge <thomas@codesourcery.com>
* gimplify.c (gimplify_omp_target_update): Elaborate 'exit data'
'finalize' handling.
2019-12-18 Tobias Burnus <tobias@codesourcery.com>
PR middle-end/86416

23
gcc/gimplify.c
View File

@ -12738,27 +12738,30 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
&& omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
OMP_CLAUSE_FINALIZE))
{
/* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
semantics apply to all mappings of this OpenACC directive. */
bool finalize_marked = false;
/* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote "finalize"
semantics. */
bool have_clause = false;
for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_FROM:
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM);
finalize_marked = true;
have_clause = true;
break;
case GOMP_MAP_RELEASE:
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
finalize_marked = true;
have_clause = true;
break;
case GOMP_MAP_POINTER:
case GOMP_MAP_TO_PSET:
/* TODO PR92929: we may see these here, but they'll always follow
one of the clauses above, and will be handled by libgomp as
one group, so no handling required here. */
gcc_assert (have_clause);
break;
default:
/* Check consistency: libgomp relies on the very first data
mapping clause being marked, so make sure we did that before
any other mapping clauses. */
gcc_assert (finalize_marked);
break;
gcc_unreachable ();
}
}
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));

5
gcc/testsuite/ChangeLog
View File

@ -1,3 +1,8 @@
2019-12-18 Thomas Schwinge <thomas@codesourcery.com>
* c-c++-common/goacc/finalize-1.c: Extend.
* gfortran.dg/goacc/finalize-1.f: Likewise.
2019-12-18 Harald Anlauf <anlauf@gmx.de>
PR fortran/70853

11
gcc/testsuite/c-c++-common/goacc/finalize-1.c
View File

@ -4,8 +4,10 @@
extern int del_r;
extern float del_f[3];
extern char *del_f_p;
extern double cpo_r[8];
extern long cpo_f;
extern char *cpo_f_p;
void f ()
{
@ -17,6 +19,10 @@ void f ()
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */
#pragma acc exit data finalize delete (del_f_p[2:5])
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) map\\(firstprivate:del_f_p \\\[pointer assign, bias: 2\\\]\\) finalize;$" 1 "original" } }
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */
#pragma acc exit data copyout (cpo_r)
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
@ -24,5 +30,8 @@ void f ()
#pragma acc exit data copyout (cpo_f) finalize
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } }
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
}
#pragma acc exit data copyout (cpo_f_p[4:10]) finalize
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\) map\\(firstprivate:cpo_f_p \\\[pointer assign, bias: 4\\\]\\);$" 1 "original" } }
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */
}

10
gcc/testsuite/gfortran.dg/goacc/finalize-1.f
View File

@ -6,8 +6,10 @@
IMPLICIT NONE
INTEGER :: del_r
REAL, DIMENSION (3) :: del_f
INTEGER (1), DIMENSION (:), ALLOCATABLE :: del_f_p
DOUBLE PRECISION, DIMENSION (8) :: cpo_r
LOGICAL :: cpo_f
INTEGER (1), DIMENSION (:), ALLOCATABLE :: cpo_f_p
!$ACC EXIT DATA DELETE (del_r)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
@ -17,6 +19,10 @@
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_r)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
@ -24,4 +30,8 @@
!$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_f\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
END SUBROUTINE f

3
libgomp/ChangeLog
View File

@ -1,5 +1,8 @@
2019-12-18 Thomas Schwinge <thomas@codesourcery.com>
* oacc-mem.c (GOACC_enter_exit_data): Simplify 'exit data'
'finalize' handling.
PR libgomp/92848
* oacc-mem.c (acc_map_data, present_create_copy)
(goacc_insert_pointer): Use 'GOMP_MAP_VARS_ENTER_DATA'.

14
libgomp/oacc-mem.c
View File

@ -1061,17 +1061,6 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
thr = goacc_thread ();
acc_dev = thr->dev;
/* Determine whether "finalize" semantics apply to all mappings of this
OpenACC directive. */
bool finalize = false;
if (mapnum > 0)
{
unsigned char kind = kinds[0] & 0xff;
if (kind == GOMP_MAP_DELETE
|| kind == GOMP_MAP_FORCE_FROM)
finalize = true;
}
/* Determine if this is an "acc enter data". */
for (i = 0; i < mapnum; ++i)
{
@ -1224,6 +1213,9 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
{
unsigned char kind = kinds[i] & 0xff;
bool finalize = (kind == GOMP_MAP_DELETE
|| kind == GOMP_MAP_FORCE_FROM);
int pointer = find_pointer (i, mapnum, kinds);
if (!pointer)