public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc r12-1363] OpenACC: Separate enter/exit data ABIs
@ 2021-06-10 11:40 Thomas Schwinge
0 siblings, 0 replies; only message in thread
From: Thomas Schwinge @ 2021-06-10 11:40 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:7aefef31365b9c3d32a0edb6ea0d3b8864d7e91a
commit r12-1363-g7aefef31365b9c3d32a0edb6ea0d3b8864d7e91a
Author: Andrew Stubbs <ams@codesourcery.com>
Date: Fri Sep 25 16:22:47 2020 +0100
OpenACC: Separate enter/exit data ABIs
Move the OpenACC enter and exit data directives from using a single builtin to
having one each. For most purposes it was easy to tell which was which, from
the clauses given, but it's overhead we can easily avoid, and there may be
future uses where that isn't possible.
gcc/
* omp-builtins.def (BUILT_IN_GOACC_ENTER_EXIT_DATA): Split into...
(BUILT_IN_GOACC_ENTER_DATA, BUILT_IN_GOACC_EXIT_DATA): ... these.
* gimple.h (enum gf_mask): Split
'GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA' into
'GF_OMP_TARGET_KIND_OACC_ENTER_DATA' and
'GF_OMP_TARGET_KIND_OACC_EXIT_DATA'.
(is_gimple_omp_oacc): Update.
* gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
* gimplify.c (gimplify_omp_target_update): Likewise.
* omp-expand.c (expand_omp_target, build_omp_regions_1)
(omp_make_gimple_edges): Likewise.
* omp-low.c (check_omp_nesting_restrictions, lower_omp_target):
Likewise.
gcc/testsuite/
* c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust patterns.
* c-c++-common/goacc/finalize-1.c: Likewise.
* c-c++-common/goacc/mdc-1.c: Likewise.
* c-c++-common/goacc/nesting-fail-1.c: Likewise.
* c-c++-common/goacc/struct-enter-exit-data-1.c: Likewise.
* gfortran.dg/goacc/attach-descriptor.f90: Likewise.
* gfortran.dg/goacc/finalize-1.f: Likewise.
* gfortran.dg/goacc/mapping-tests-3.f90: Likewise.
libgomp/
* libgomp.map (GOACC_2.0.2): New symbol version.
* libgomp_g.h (GOACC_enter_data, GOACC_exit_data) New prototypes.
* oacc-mem.c (GOACC_enter_data, GOACC_exit_data) New functions.
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Diff:
---
gcc/gimple-pretty-print.c | 7 ++++--
gcc/gimple.h | 18 ++++++++-------
gcc/gimplify.c | 5 ++++-
gcc/omp-builtins.def | 5 ++++-
gcc/omp-expand.c | 19 +++++++++++-----
gcc/omp-low.c | 9 +++++---
.../c-c++-common/goacc-gomp/nesting-fail-1.c | 4 ++--
gcc/testsuite/c-c++-common/goacc/finalize-1.c | 12 +++++-----
gcc/testsuite/c-c++-common/goacc/mdc-1.c | 18 +++++++--------
gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c | 8 +++----
.../c-c++-common/goacc/struct-enter-exit-data-1.c | 4 ++--
.../gfortran.dg/goacc/attach-descriptor.f90 | 6 ++---
gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 12 +++++-----
.../gfortran.dg/goacc/mapping-tests-3.f90 | 2 +-
libgomp/libgomp.map | 6 +++++
libgomp/libgomp_g.h | 4 ++++
libgomp/oacc-mem.c | 26 ++++++++++++++++++++++
17 files changed, 111 insertions(+), 54 deletions(-)
diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c
index c9c0a66ee30..8be40416dd2 100644
--- a/gcc/gimple-pretty-print.c
+++ b/gcc/gimple-pretty-print.c
@@ -1697,8 +1697,11 @@ dump_gimple_omp_target (pretty_printer *buffer, const gomp_target *gs,
case GF_OMP_TARGET_KIND_OACC_UPDATE:
kind = " oacc_update";
break;
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
- kind = " oacc_enter_exit_data";
+ case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
+ kind = " oacc_enter_data";
+ break;
+ case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
+ kind = " oacc_exit_data";
break;
case GF_OMP_TARGET_KIND_OACC_DECLARE:
kind = " oacc_declare";
diff --git a/gcc/gimple.h b/gcc/gimple.h
index 91b92b4a4d1..e7dc2a45a13 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -161,7 +161,7 @@ enum gf_mask {
GF_OMP_FOR_KIND_SIMD = 5,
GF_OMP_FOR_COMBINED = 1 << 3,
GF_OMP_FOR_COMBINED_INTO = 1 << 4,
- GF_OMP_TARGET_KIND_MASK = (1 << 4) - 1,
+ GF_OMP_TARGET_KIND_MASK = (1 << 5) - 1,
GF_OMP_TARGET_KIND_REGION = 0,
GF_OMP_TARGET_KIND_DATA = 1,
GF_OMP_TARGET_KIND_UPDATE = 2,
@@ -172,18 +172,19 @@ enum gf_mask {
GF_OMP_TARGET_KIND_OACC_SERIAL = 7,
GF_OMP_TARGET_KIND_OACC_DATA = 8,
GF_OMP_TARGET_KIND_OACC_UPDATE = 9,
- GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 10,
- GF_OMP_TARGET_KIND_OACC_DECLARE = 11,
- GF_OMP_TARGET_KIND_OACC_HOST_DATA = 12,
+ GF_OMP_TARGET_KIND_OACC_ENTER_DATA = 10,
+ GF_OMP_TARGET_KIND_OACC_EXIT_DATA = 11,
+ GF_OMP_TARGET_KIND_OACC_DECLARE = 12,
+ GF_OMP_TARGET_KIND_OACC_HOST_DATA = 13,
/* A 'GF_OMP_TARGET_KIND_OACC_PARALLEL' representing an OpenACC 'kernels'
decomposed part, parallelized. */
- GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED = 13,
+ GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED = 14,
/* A 'GF_OMP_TARGET_KIND_OACC_PARALLEL' representing an OpenACC 'kernels'
decomposed part, "gang-single". */
- GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE = 14,
+ GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE = 15,
/* A 'GF_OMP_TARGET_KIND_OACC_DATA' representing an OpenACC 'kernels'
decomposed parts' 'data' construct. */
- GF_OMP_TARGET_KIND_OACC_DATA_KERNELS = 15,
+ GF_OMP_TARGET_KIND_OACC_DATA_KERNELS = 16,
GF_OMP_TEAMS_HOST = 1 << 0,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
@@ -6525,7 +6526,8 @@ is_gimple_omp_oacc (const gimple *stmt)
case GF_OMP_TARGET_KIND_OACC_SERIAL:
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
+ case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_DECLARE:
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED:
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 39f5b973d18..c96d611b115 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -13383,8 +13383,11 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
case OACC_ENTER_DATA:
+ kind = GF_OMP_TARGET_KIND_OACC_ENTER_DATA;
+ ort = ORT_ACC;
+ break;
case OACC_EXIT_DATA:
- kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
+ kind = GF_OMP_TARGET_KIND_OACC_EXIT_DATA;
ort = ORT_ACC;
break;
case OACC_UPDATE:
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index cfbf1e67b8e..97964f866ec 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -35,7 +35,10 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
BT_FN_VOID, ATTR_NOTHROW_LIST)
-DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_EXIT_DATA, "GOACC_enter_exit_data",
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_DATA, "GOACC_enter_data",
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
+ ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_EXIT_DATA, "GOACC_exit_data",
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel_keyed",
diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 0f843bad79a..f8b1558b1d1 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -9290,7 +9290,8 @@ expand_omp_target (struct omp_region *region)
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_SERIAL:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
+ case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_DECLARE:
case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED:
case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE:
@@ -9574,8 +9575,11 @@ expand_omp_target (struct omp_region *region)
case GF_OMP_TARGET_KIND_OACC_UPDATE:
start_ix = BUILT_IN_GOACC_UPDATE;
break;
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
- start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
+ start_ix = BUILT_IN_GOACC_ENTER_DATA;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
+ start_ix = BUILT_IN_GOACC_EXIT_DATA;
break;
case GF_OMP_TARGET_KIND_OACC_DECLARE:
start_ix = BUILT_IN_GOACC_DECLARE;
@@ -9773,7 +9777,8 @@ expand_omp_target (struct omp_region *region)
oacc_set_fn_attrib (child_fn, clauses, &args);
tagging = true;
/* FALLTHRU */
- case BUILT_IN_GOACC_ENTER_EXIT_DATA:
+ case BUILT_IN_GOACC_ENTER_DATA:
+ case BUILT_IN_GOACC_EXIT_DATA:
case BUILT_IN_GOACC_UPDATE:
{
tree t_async = NULL_TREE;
@@ -10042,7 +10047,8 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
+ case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_DECLARE:
/* ..., other than for those stand-alone directives... */
region = NULL;
@@ -10299,7 +10305,8 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region,
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
+ case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_DECLARE:
cur_region = cur_region->outer;
break;
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 2d5cdf671eb..2cc2a186080 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -3646,8 +3646,10 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
case GF_OMP_TARGET_KIND_OACC_SERIAL: stmt_name = "serial"; break;
case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name = "data"; break;
case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break;
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
- stmt_name = "enter/exit data"; break;
+ case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
+ stmt_name = "enter data"; break;
+ case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
+ stmt_name = "exit data"; break;
case GF_OMP_TARGET_KIND_OACC_DECLARE: stmt_name = "declare"; break;
case GF_OMP_TARGET_KIND_OACC_HOST_DATA: stmt_name = "host_data";
break;
@@ -12194,7 +12196,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_SERIAL:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
- case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
+ case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
+ case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_DECLARE:
case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED:
case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE:
diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 1a3324200e2..ddbd247342f 100644
--- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -162,8 +162,8 @@ f_omp (void)
#pragma acc data /* { dg-error "OpenACC .data. construct inside of OpenMP .target. region" } */
;
#pragma acc update host(i) /* { dg-error "OpenACC .update. construct inside of OpenMP .target. region" } */
-#pragma acc enter data copyin(i) /* { dg-error "OpenACC .enter/exit data. construct inside of OpenMP .target. region" } */
-#pragma acc exit data delete(i) /* { dg-error "OpenACC .enter/exit data. construct inside of OpenMP .target. region" } */
+#pragma acc enter data copyin(i) /* { dg-error "OpenACC .enter data. construct inside of OpenMP .target. region" } */
+#pragma acc exit data delete(i) /* { dg-error "OpenACC .exit data. construct inside of OpenMP .target. region" } */
#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
for (i = 0; i < 2; ++i)
;
diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
index 3d64b2e7cb3..54bf1b76a1b 100644
--- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
@@ -13,25 +13,25 @@ void f ()
{
#pragma acc exit data delete (del_r)
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
- { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
#pragma acc exit data finalize delete (del_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" } } */
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_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" } } */
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_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" } } */
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
#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" } } */
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_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" } } */
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */
}
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
index c9ab7c24074..c2b8dc6c880 100644
--- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -43,15 +43,15 @@ t1 ()
}
}
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:s .len: 32.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.attach:a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:a .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.release:a .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
index 93a911120d4..5cfb327f4b6 100644
--- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -14,8 +14,8 @@ f_acc_parallel (void)
#pragma acc data /* { dg-error ".data. construct inside of .parallel. region" } */
;
#pragma acc update host(i) /* { dg-error ".update. construct inside of .parallel. region" } */
-#pragma acc enter data copyin(i) /* { dg-error ".enter/exit data. construct inside of .parallel. region" } */
-#pragma acc exit data delete(i) /* { dg-error ".enter/exit data. construct inside of .parallel. region" } */
+#pragma acc enter data copyin(i) /* { dg-error ".enter data. construct inside of .parallel. region" } */
+#pragma acc exit data delete(i) /* { dg-error ".exit data. construct inside of .parallel. region" } */
}
}
@@ -33,8 +33,8 @@ f_acc_kernels (void)
#pragma acc data /* { dg-error ".data. construct inside of .kernels. region" } */
;
#pragma acc update host(i) /* { dg-error ".update. construct inside of .kernels. region" } */
-#pragma acc enter data copyin(i) /* { dg-error ".enter/exit data. construct inside of .kernels. region" } */
-#pragma acc exit data delete(i) /* { dg-error ".enter/exit data. construct inside of .kernels. region" } */
+#pragma acc enter data copyin(i) /* { dg-error ".enter data. construct inside of .kernels. region" } */
+#pragma acc exit data delete(i) /* { dg-error ".exit data. construct inside of .kernels. region" } */
}
}
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
index df405e448b2..9e5d3f2c9d2 100644
--- 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
@@ -20,8 +20,8 @@ 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 } } */
+ /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_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 } } */
+ /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_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 } } */
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
index 373bdcb2114..8c2ee4a5cca 100644
--- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
@@ -12,11 +12,11 @@ program att
!$acc enter data attach(myvar%arr2, myptr)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
!$acc exit data detach(myvar%arr2, myptr)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
! Test valid usage and processing of the finalize clause.
!$acc exit data detach(myvar%arr2, myptr) finalize
@@ -24,6 +24,6 @@ program att
! For array-descriptor detaches, we no longer generate a "release" mapping
! for the pointed-to data for gimplify.c to turn into "delete". Make sure
! the mapping still isn't there.
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
end program att
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index a7788580819..b706b385aeb 100644
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -13,25 +13,25 @@
!$ACC EXIT DATA DELETE (del_r)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
!$ACC EXIT DATA FINALIZE DELETE (del_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" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_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: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) 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" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_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" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
!$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" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_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: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) 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" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_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
diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
index 890ca781967..662104f8aaa 100644
--- a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90
@@ -11,5 +11,5 @@ subroutine foo
type(two) x
!$acc enter data copyin(x%A)
-! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_enter_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } }
end
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 4ad190a52af..8ea27b5565f 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -541,6 +541,12 @@ GOACC_2.0.1 {
GOACC_parallel_keyed;
} GOACC_2.0;
+GOACC_2.0.2 {
+ global:
+ GOACC_enter_data;
+ GOACC_exit_data;
+} GOACC_2.0.1;
+
GOMP_PLUGIN_1.0 {
global:
GOMP_PLUGIN_malloc;
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index b66b6978202..f890a204fb8 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -370,6 +370,10 @@ extern void GOACC_wait (int, int, ...);
extern void GOACC_enter_exit_data (int, size_t, void **, size_t *,
unsigned short *, int, int, ...);
+extern void GOACC_enter_data (int, size_t, void **, size_t *,
+ unsigned short *, int, int, ...);
+extern void GOACC_exit_data (int, size_t, void **, size_t *,
+ unsigned short *, int, int, ...);
extern void GOACC_declare (int, size_t, void **, size_t *, unsigned short *);
/* oacc-parallel.c */
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index f6173b91fdd..123fe154089 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1422,6 +1422,8 @@ goacc_enter_exit_data_internal (int flags_m, size_t mapnum, void **hostaddrs,
}
}
+/* Legacy entry point (GCC 11 and earlier). */
+
void
GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
size_t *sizes, unsigned short *kinds, int async,
@@ -1468,6 +1470,30 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
va_end (ap);
}
+void
+GOACC_enter_data (int flags_m, size_t mapnum, void **hostaddrs,
+ size_t *sizes, unsigned short *kinds, int async,
+ int num_waits, ...)
+{
+ va_list ap;
+ va_start (ap, num_waits);
+ goacc_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds,
+ true, async, num_waits, &ap);
+ va_end (ap);
+}
+
+void
+GOACC_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
+ size_t *sizes, unsigned short *kinds, int async,
+ int num_waits, ...)
+{
+ va_list ap;
+ va_start (ap, num_waits);
+ goacc_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds,
+ false, async, num_waits, &ap);
+ va_end (ap);
+}
+
void
GOACC_declare (int flags_m, size_t mapnum,
void **hostaddrs, size_t *sizes, unsigned short *kinds)
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-06-10 11:40 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-06-10 11:40 [gcc r12-1363] OpenACC: Separate enter/exit data ABIs Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).