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 --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c
index c9c0a66ee308a034cc82b4976c15a0e20e43a4b7..8be40416dd2b3c74e2eaae8c678bfc17ff79f872 100644 (file)
--- 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 91b92b4a4d172d595fdd1b26f8a5c317f13775f4..e7dc2a45a13e62fd9b2917772848892774eed91b 100644 (file)
--- 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 39f5b973d18132c7f056bac4fbe0bee7f3afe43f..c96d611b11583dec5b20224e9e572e3628924958 100644 (file)
--- 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 cfbf1e67b8e5f807e818e3aa9151da3a2be7dcb6..97964f866ec05581e3b4220c19cc5f532469e9be 100644 (file)
--- 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 0f843bad79a29864df8ba25398d43a71033d7491..f8b1558b1d1ba9a79eebfb8ebd6cd620ddd07807 100644 (file)
--- 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 2d5cdf671ebd365d6276f20fbd71cd0393ef64bc..2cc2a1860806e6d5efb996b20817320f2cbfc92d 100644 (file)
--- 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 1a3324200e2c4989d5189cafc35df7bf38d8aa0b..ddbd247342f026cf23831934ed51d28635397c34 100644 (file)
--- 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 3d64b2e7cb32f75cc5c0077d35902d1edd8be0c6..54bf1b76a1b8cb956240684cfc8b67b9a7426ae8 100644 (file)
--- 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 c9ab7c240746934b2efa2325b563ddf633ff8e7d..c2b8dc6c88088e703c95f8ad0d8aa024280c3c65 100644 (file)
--- 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 93a911120d486dd94e1ba1fd96e208a17a6960a7..5cfb327f4b60951e785dd35f39409137cbcd0b09 100644 (file)
--- 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 df405e448b28933a81341e1bb41741fe54c32b25..9e5d3f2c9d207a7e0123be6c251cc64ee88e685a 100644 (file)
--- 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 373bdcb21145a69566e8fa5306c736a463343a33..8c2ee4a5cca449c8577ff8f4713dee46265637d2 100644 (file)
--- 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 a77885808199357530440abfd877261ae96b25e1..b706b385aebcdc87c0db053990b7863ca5036485 100644 (file)
--- 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 890ca781967d1fcf408186a58593cd7aefe2e8c4..662104f8aaaba0b762ee702d95d191d975c32e5d 100644 (file)
--- 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 4ad190a52af099e55f71ac639e03098b607a076b..8ea27b5565ff61500f24b3cf67b1f8f3e6380456 100644 (file)
--- 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 b66b69782029638f478c28b831d54c8990187848..f890a204fb8ce82cc0120be09b740f66123201fb 100644 (file)
--- 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 f6173b91fddaa3396bfe3781edd6acf2f7860855..123fe1540896d488b7b060f75afe526a2bccbe22 100644 (file)
--- 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)