From 1dce4fd66af3e6c94eef205d9dac42800f712f37 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 20 May 2020 10:56:55 +0200 Subject: [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) : Explain special handling. Co-Authored-By: Julian Brown --- gcc/gimplify.c | 3 ++- .../c-c++-common/goacc/struct-enter-exit-data-1.c | 27 ++++++++++++++++++++++ 2 files changed, 29 insertions(+), 1 deletion(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c (limited to 'gcc') diff --git a/gcc/gimplify.c b/gcc/gimplify.c index cb08b26..e14932f 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -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 diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c new file mode 100644 index 0000000..df405e4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c @@ -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 } } */ +} -- cgit v1.1