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

[OpenACC 'exit data'] Strip 'GOMP_MAP_STRUCT' mappings 

These are not itself necessary for OpenACC 'exit data' directives, and are
skipped over (now) in libgomp.  We might as well not emit them to start with,
in line with the equivalent OpenMP directive.  We keep the no-op handling in
libgomp for the reason of backward compatibility.

	gcc/
	* gimplify.c (gimplify_adjust_omp_clauses): Remove
	'GOMP_MAP_STRUCT' mapping from OpenACC 'exit data' directives.
	gcc/testsuite/
	* c-c++-common/goacc/struct-enter-exit-data-1.c: New file.
	libgomp/
	* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>: Explain
	special handling.

Co-Authored-By: Julian Brown <julian@codesourcery.com>
This commit is contained in:
Thomas Schwinge 2020-05-20 10:56:55 +02:00
parent 1809628fcf
commit 1afc467256
3 changed files with 32 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
gcc/gimplify.c
View File

@ -10396,7 +10396,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
}
else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
&& code == OMP_TARGET_EXIT_DATA)
&& (code == OMP_TARGET_EXIT_DATA
|| code == OACC_EXIT_DATA))
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST

27
gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c Normal file
View File

@ -0,0 +1,27 @@
/* Check 'GOMP_MAP_STRUCT' mapping, and in particular that it gets removed from
OpenACC 'exit data' directives. */
/* { dg-additional-options "-fdump-tree-gimple" } */
struct str {
int a;
int *b;
int *c;
int d;
int *e;
int f;
};
#define N 1024
void
test (int *b, int *c, int *e)
{
struct str s = { .a = 0, .b = b, .c = c, .d = 0, .e = e, .f = 0 };
#pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
/* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
#pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
/* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
}

5
libgomp/oacc-mem.c
View File

@ -1181,8 +1181,9 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
case GOMP_MAP_STRUCT:
/* Skip the 'GOMP_MAP_STRUCT' itself, and use the regular processing
for all its entries. TODO: don't generate these no-op
'GOMP_MAP_STRUCT's. */
for all its entries. This special handling exists for GCC 10.1
compatibility; afterwards, we're not generating these no-op
'GOMP_MAP_STRUCT's anymore. */
break;
default: