OpenMP: Cleanups related to the 'present' modifier

Reduce number of enum values passed to libgomp as
GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} have the same semantic as
GOMP_MAP_FORCE_PRESENT (i.e. abort if not present, otherwise ignore);
that's different to GOMP_MAP_ALWAYS_PRESENT_{TO,TOFROM,FROM} which also
abort if not present but copy data when present. This is is a follow-up to
the commit r14-1579-g4ede915d5dde93 done 6 days ago.

Additionally, the commit improves a libgomp run-time and a C/C++ compile-time
error wording and extends testcases a tiny bit.

gcc/c/ChangeLog:

	* c-parser.cc (c_parser_omp_clause_map): Reword error message for
	clearness especially with 'omp target (enter/exit) data.'

gcc/cp/ChangeLog:

	* parser.cc (cp_parser_omp_clause_map): Reword error message for
	clearness especially with 'omp target (enter/exit) data.'
	* semantics.cc (handle_omp_array_sections): Handle
	GOMP_MAP_{ALWAYS_,}PRESENT_{TO,TOFROM,FROM,ALLOC} enum values.

gcc/ChangeLog:

	* gimplify.cc (gimplify_adjust_omp_clauses_1): Use
	GOMP_MAP_FORCE_PRESENT for 'present alloc' implicit mapping.
	(gimplify_adjust_omp_clauses): Change
	GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} to the equivalent
	GOMP_MAP_FORCE_PRESENT.
	* omp-low.cc (lower_omp_target): Remove handling of no-longer valid
	GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC}; update map kinds used for
	to/from clauses with present modifier.

include/ChangeLog:

	* gomp-constants.h (enum gomp_map_kind): Change the enum values
	GOMP_MAP_PRESENT_{TO,TOFROM,FROM,ALLOC} to be compiler only.
	(GOMP_MAP_PRESENT_P): Update to include also GOMP_MAP_FORCE_PRESENT.

libgomp/ChangeLog:

	* target.c (gomp_to_device_kind_p, gomp_map_vars_internal): Replace
	GOMP_MAP_PRESENT_{FROM,TO,TOFROM,ACLLOC} by GOMP_MAP_FORCE_PRESENT.
	(gomp_map_vars_internal, gomp_update): Likewise; unify and improve
	error message.
	* testsuite/libgomp.c-c++-common/target-present-2.c: Update for
	changed error message.
	* testsuite/libgomp.fortran/target-present-1.f90: Likewise.
	* testsuite/libgomp.fortran/target-present-2.f90: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/present-1.c: Likewise.
	* testsuite/libgomp.c-c++-common/target-present-1.c: Likewise and
	extend testcase to check that data is copied when needed.
	* testsuite/libgomp.c-c++-common/target-present-3.c: Likewise.
	* testsuite/libgomp.fortran/target-present-3.f90: Likewise.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/defaultmap-4.c: Update scan-tree-dump.
	* c-c++-common/gomp/map-9.c: Likewise.
	* gfortran.dg/gomp/defaultmap-8.f90: Likewise.
	* gfortran.dg/gomp/map-11.f90: Likewise.
	* gfortran.dg/gomp/target-update-1.f90: Likewise.
	* gfortran.dg/gomp/map-12.f90: Likewise; also check original dump.
	* c-c++-common/gomp/map-6.c: Update dg-error and also check
	clause error with 'target (enter/exit) data'.

diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 72f6fbae6a69349a3eec06f99b717526eb6e9cd6..61487cc7481b16d5610bfad365b64952efdb1050 100644 (file)
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -17160,9 +17160,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
        }
       else
        {
-         c_parser_error (parser, "%<#pragma omp target%> with "
-                                 "modifier other than %<always%>, %<close%> "
-                                 "or %<present%> on %<map%> clause");
+         c_parser_error (parser, "%<map%> clause with map-type modifier other "
+                                 "than %<always%>, %<close%> or %<present%>");
          parens.skip_until_found_close (parser);
          return list;
        }

diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 09cba71343741b1136fb2c7c9f4fca1657ed8113..f5f9f5a45104ab36a66068f25437e4056da88a48 100644 (file)
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -40625,9 +40625,8 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
        }
       else
        {
-         cp_parser_error (parser, "%<#pragma omp target%> with "
-                                  "modifier other than %<always%>, %<close%> "
-                                  "or %<present%> on %<map%> clause");
+         cp_parser_error (parser, "%<map%> clause with map-type modifier other"
+                                  " than %<always%>, %<close%> or %<present%>");
          cp_parser_skip_to_closing_parenthesis (parser,
                                                 /*recovering=*/true,
                                                 /*or_comma=*/false,

diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index a2e74a5d2c77c6158ab6a944bcff6c052417ddc5..8fb47fd179eb2af2e82bf31d188023e9b9d41de9 100644 (file)
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5819,6 +5819,13 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
            case GOMP_MAP_ALWAYS_TO:
            case GOMP_MAP_ALWAYS_FROM:
            case GOMP_MAP_ALWAYS_TOFROM:
+           case GOMP_MAP_PRESENT_ALLOC:
+           case GOMP_MAP_PRESENT_TO:
+           case GOMP_MAP_PRESENT_FROM:
+           case GOMP_MAP_PRESENT_TOFROM:
+           case GOMP_MAP_ALWAYS_PRESENT_TO:
+           case GOMP_MAP_ALWAYS_PRESENT_FROM:
+           case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
            case GOMP_MAP_RELEASE:
            case GOMP_MAP_DELETE:
            case GOMP_MAP_FORCE_TO:

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 91640deecea25ea3388284e3977a45f200afb883..0e24b915b8fe4a3886221f4080e94a56cb938478 100644 (file)
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -12479,7 +12479,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
          kind = GOMP_MAP_FORCE_PRESENT;
          break;
        case GOVD_MAP_FORCE_PRESENT | GOVD_MAP_ALLOC_ONLY:
-         kind = GOMP_MAP_PRESENT_ALLOC;
+         kind = GOMP_MAP_FORCE_PRESENT;
          break;
        default:
          gcc_unreachable ();
@@ -12797,6 +12797,17 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
          break;
 
        case OMP_CLAUSE_MAP:
+         switch (OMP_CLAUSE_MAP_KIND (c))
+           {
+           case GOMP_MAP_PRESENT_ALLOC:
+           case GOMP_MAP_PRESENT_TO:
+           case GOMP_MAP_PRESENT_FROM:
+           case GOMP_MAP_PRESENT_TOFROM:
+             OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_PRESENT);
+             break;
+           default:
+             break;
+           }
          if (code == OMP_TARGET_EXIT_DATA
              && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
            {

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 1857b5b960f305afb12f998f13db6e53cd29013a..b882df048ef20df77c785b35ee8f2fffe3224e2f 100644 (file)
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -12800,10 +12800,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
          case GOMP_MAP_ALWAYS_TO:
          case GOMP_MAP_ALWAYS_FROM:
          case GOMP_MAP_ALWAYS_TOFROM:
-         case GOMP_MAP_PRESENT_ALLOC:
-         case GOMP_MAP_PRESENT_FROM:
-         case GOMP_MAP_PRESENT_TO:
-         case GOMP_MAP_PRESENT_TOFROM:
+         case GOMP_MAP_FORCE_PRESENT:
          case GOMP_MAP_ALWAYS_PRESENT_FROM:
          case GOMP_MAP_ALWAYS_PRESENT_TO:
          case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
@@ -12822,7 +12819,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
          case GOMP_MAP_FORCE_TO:
          case GOMP_MAP_FORCE_FROM:
          case GOMP_MAP_FORCE_TOFROM:
-         case GOMP_MAP_FORCE_PRESENT:
          case GOMP_MAP_FORCE_DEVICEPTR:
          case GOMP_MAP_DEVICE_RESIDENT:
          case GOMP_MAP_LINK:
@@ -13349,10 +13345,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                    case GOMP_MAP_ALWAYS_TO:
                    case GOMP_MAP_ALWAYS_FROM:
                    case GOMP_MAP_ALWAYS_TOFROM:
-                   case GOMP_MAP_PRESENT_ALLOC:
-                   case GOMP_MAP_PRESENT_TO:
-                   case GOMP_MAP_PRESENT_FROM:
-                   case GOMP_MAP_PRESENT_TOFROM:
                    case GOMP_MAP_ALWAYS_PRESENT_TO:
                    case GOMP_MAP_ALWAYS_PRESENT_FROM:
                    case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
@@ -13397,13 +13389,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
              case OMP_CLAUSE_TO:
                tkind
                  = (OMP_CLAUSE_MOTION_PRESENT (c)
-                    ? GOMP_MAP_PRESENT_TO : GOMP_MAP_TO);
+                    ? GOMP_MAP_ALWAYS_PRESENT_TO : GOMP_MAP_TO);
                tkind_zero = tkind;
                break;
              case OMP_CLAUSE_FROM:
                tkind
                  = (OMP_CLAUSE_MOTION_PRESENT (c)
-                    ? GOMP_MAP_PRESENT_FROM : GOMP_MAP_FROM);
+                    ? GOMP_MAP_ALWAYS_PRESENT_FROM : GOMP_MAP_FROM);
                tkind_zero = tkind;
                break;
              default:

diff --git a/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c
index 1afff7ea38f552cb2c03f6641785372abc0e4cac..b84f89b0c7c792a3cfee37fb4b98eb737232b3a0 100644 (file)
--- a/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c
+++ b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c
@@ -20,5 +20,5 @@ foo (void)
       c[i] = a[i] + b[i];
 }
 
-/* { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\) map\\(present,alloc:c \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" "gimple" } } */
-/* { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\) map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\) map\\(force_present:c \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */

diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
index 8c5b7f7cca101684cb9ca8f8861b4cb7e794c129..5152d9d7c21a67087b0160542d6f113f9e1699cc 100644 (file)
--- a/gcc/testsuite/c-c++-common/gomp/map-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -13,12 +13,22 @@ foo (void)
   #pragma omp target map (to:a)
   ;
 
-  #pragma omp target map (a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'present'" } */
+  #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
   ;
 
-  #pragma omp target map (close, a to: b) /* { dg-error "'#pragma omp target' with modifier other than 'always', 'close' or 'present'" } */
+  #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
   ;
 
+  #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+  ;
+
+  #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+  ;
+
+  #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+  ;
+
+
   #pragma omp target map (close a) /* { dg-error "'close' undeclared" "" { target c } } */ 
   /* { dg-error "'close' has not been declared" "" { target c++ } .-1 } */ 
   /* { dg-error "expected '\\)' before 'a'" "" { target *-*-* } .-2 } */

diff --git a/gcc/testsuite/c-c++-common/gomp/map-9.c b/gcc/testsuite/c-c++-common/gomp/map-9.c
index 4b4bd6d2aa32be66fe307aa3c4b8b2230808f836..fcf3125ceeeddc8b53f7b8ab769e8b9c3a16608b 100644 (file)
--- a/gcc/testsuite/c-c++-common/gomp/map-9.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-9.c
@@ -25,8 +25,8 @@ foo (void)
       c[i] = a[i] + b[i];
 }
 
-/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
-/* { dg-final { scan-tree-dump "pragma omp target data map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
-/* { dg-final { scan-tree-dump "pragma omp target.*map\\(present,from:c \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target data map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\) map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target.*map\\(force_present:c \\\[len: \[0-9\]+\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
 /* { dg-final { scan-tree-dump "pragma omp target exit data map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
-/* { dg-final { scan-tree-dump "pragma omp target.*map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\) map\\(to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target.*map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\) map\\(to:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */

diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90
index 669a623f7460e350601c0b38453b7d820411d592..e26d1e004b1a711f4215be47d63ec1b5ec1a1002 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90
@@ -22,5 +22,5 @@ program main
   !$omp end target
 end program
   
-! { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\).*map\\(present,alloc:c \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" "gimple" } }
-! { dg-final { scan-tree-dump "pragma omp target.*map\\(present,alloc:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(present,alloc:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\) defaultmap\\(present:aggregate\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target.*defaultmap\\(present:aggregate\\).*map\\(force_present:c \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target.*map\\(force_present:b \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(force_present:a \\\[len: \[0-9\]+\\\]\\\[implicit\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\) defaultmap\\(present:aggregate\\)" "gimple" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/map-11.f90 b/gcc/testsuite/gfortran.dg/gomp/map-11.f90
index 9eb956f1aa1c3bf3a67155519052f6dc88479fda..7ef9d46f2f821c746c31efbc29e58f71de842e4c 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/map-11.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-11.f90
@@ -27,8 +27,8 @@ program main
   !$omp end target
 end program
 
-! { dg-final { scan-tree-dump "pragma omp target enter data map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
-! { dg-final { scan-tree-dump "pragma omp target data map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
-! { dg-final { scan-tree-dump "pragma omp target.*map\\(present,to:a \\\[len: \[0-9\]+\\\]\\) map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target enter data map\\(force_present:a \\\[len: \[0-9\]+\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target data map\\(force_present:a \\\[len: \[0-9\]+\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target.*map\\(force_present:a \\\[len: \[0-9\]+\\\]\\) map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(force_present:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
 ! { dg-final { scan-tree-dump "pragma omp target exit data map\\(always,present,from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
-! { dg-final { scan-tree-dump "pragma omp target.*map\\(present,to:b \\\[len: \[0-9\]+\\\]\\) map\\(to:a \\\[len: \[0-9\]+\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "pragma omp target.*map\\(force_present:b \\\[len: \[0-9\]+\\\]\\) map\\(to:a \\\[len: \[0-9\]+\\\]\\) map\\(from:c \\\[len: \[0-9\]+\\\]\\)" "gimple" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/map-12.f90 b/gcc/testsuite/gfortran.dg/gomp/map-12.f90
index 74bd01f5f5a2b8c8f957cfe7de1fc3ec421e56b7..ac9a0f8aae0498712f6f155215775ab696f55c9c 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/map-12.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-12.f90
@@ -1,4 +1,4 @@
-! { dg-additional-options "-fdump-tree-omplower" }
+! { dg-additional-options "-fdump-tree-omplower -fdump-tree-original" }
 
 subroutine foo
   implicit none
@@ -40,28 +40,29 @@ subroutine foo
   !$omp target exit data map(present,always,close,delete: a) map(close,present,always,release:b) map(always close present,from:b1)
 end subroutine
 
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,tofrom:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(present,tofrom:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,present,tofrom:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(alloc:a\\) map\\(to:b\\) map\\(to:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(alloc:a\\) map\\(always,to:b\\) map\\(always,to:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a\\) map\\(present,to:b\\) map\\(present,to:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a\\) map\\(always,present,to:b\\) map\\(always,present,to:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(delete:a\\) map\\(release:b\\) map\\(from:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(delete:a\\) map\\(release:b\\) map\\(always,from:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(delete:a\\) map\\(release:b\\) map\\(present,from:b1\\)\[\r\n\]" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(delete:a\\) map\\(release:b\\) map\\(always,present,from:b1\\)\[\r\n\]" 2 "original" } }
+
 
 ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(present,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(present,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(force_present:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,present,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,present,tofrom:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(to:b \\\[len: 4\\\]\\) map\\(to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(to:b \\\[len: 4\\\]\\) map\\(to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(always,to:b \\\[len: 4\\\]\\) map\\(always,to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(always,to:b \\\[len: 4\\\]\\) map\\(always,to:b1 \\\[len: 4\\\]\\) map\\(alloc:a \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a \\\[len: 4\\\]\\) map\\(present,to:b \\\[len: 4\\\]\\) map\\(present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a \\\[len: 4\\\]\\) map\\(present,to:b \\\[len: 4\\\]\\) map\\(present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a \\\[len: 4\\\]\\) map\\(always,present,to:b \\\[len: 4\\\]\\) map\\(always,present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,alloc:a \\\[len: 4\\\]\\) map\\(always,present,to:b \\\[len: 4\\\]\\) map\\(always,present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(force_present:a \\\[len: 4\\\]\\) map\\(force_present:b \\\[len: 4\\\]\\) map\\(force_present:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(force_present:a \\\[len: 4\\\]\\) map\\(always,present,to:b \\\[len: 4\\\]\\) map\\(always,present,to:b1 \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(present,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(present,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,present,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(force_present:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }
 ! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,present,from:b1 \\\[len: 4\\\]\\) map\\(delete:a \\\[len: 4\\\]\\) map\\(release:b \\\[len: 4\\\]\\)\[\r\n\]" 2 "omplower" } }

diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90
index f99bffe2e0b0b14061825b07946acdfcf81411ed..a9db2f1a39f65ae6fb02fb3881c51a51ff5c620e 100644 (file)
--- a/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90
@@ -10,4 +10,4 @@ program main
   !$omp target update to(c) to(present: a) from(d) from(present: b) to(e)
 end program
 
-! { dg-final { scan-tree-dump "#pragma omp target update to\\(c \\\[len: \[0-9\]+\\\]\\) to\\(present:a \\\[len: \[0-9\]+\\\]\\) to\\(e \\\[len: \[0-9\]+\\\]\\) from\\(present:b \\\[len: \[0-9\]+\\\]\\) from\\(d \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
+! { dg-final { scan-tree-dump "#pragma omp target update to\\(c \\\[len: \[0-9\]+\\\]\\) to\\(present:a \\\[len: \[0-9\]+\\\]\\) to\\(e \\\[len: \[0-9\]+\\\]\\) from\\(d \\\[len: \[0-9\]+\\\]\\) from\\(present:b \\\[len: \[0-9\]+\\\]\\)" "gimple" } }

diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 49b7dd86ff5b09cc069411094b39f21119fe53bf..8d4e8e813031884c9cdf285050db55a24db0e9dc 100644 (file)
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -136,14 +136,6 @@ enum gomp_map_kind
        device.  */
     GOMP_MAP_ALWAYS_TOFROM =           (GOMP_MAP_FLAG_SPECIAL_2
                                         | GOMP_MAP_TOFROM),
-    /* Must already be present.  */
-    GOMP_MAP_PRESENT_ALLOC =           (GOMP_MAP_FLAG_PRESENT | GOMP_MAP_ALLOC),
-    /* Must already be present, copy to device.  */
-    GOMP_MAP_PRESENT_TO =              (GOMP_MAP_FLAG_PRESENT | GOMP_MAP_TO),
-    /* Must already be present, copy from device.  */
-    GOMP_MAP_PRESENT_FROM =            (GOMP_MAP_FLAG_PRESENT | GOMP_MAP_FROM),
-    /* Must already be present, copy to and from device.  */
-    GOMP_MAP_PRESENT_TOFROM =          (GOMP_MAP_FLAG_PRESENT | GOMP_MAP_TOFROM),
     /* Must already be present, unconditionally copy to device.  */
     GOMP_MAP_ALWAYS_PRESENT_TO =       (GOMP_MAP_FLAG_ALWAYS_PRESENT
                                         | GOMP_MAP_TO),
@@ -205,7 +197,13 @@ enum gomp_map_kind
     /* An attach or detach operation.  Rewritten to the appropriate type during
        gimplification, depending on directive (i.e. "enter data" or
        parallel/kernels region vs. "exit data").  */
-    GOMP_MAP_ATTACH_DETACH =           (GOMP_MAP_LAST | 3)
+    GOMP_MAP_ATTACH_DETACH =           (GOMP_MAP_LAST | 3),
+    /* Must already be present - all of following map to GOMP_MAP_FORCE_PRESENT
+       as no data transfer is needed.  */
+    GOMP_MAP_PRESENT_ALLOC =           (GOMP_MAP_LAST | 4),
+    GOMP_MAP_PRESENT_TO =              (GOMP_MAP_LAST | 5),
+    GOMP_MAP_PRESENT_FROM =            (GOMP_MAP_LAST | 6),
+    GOMP_MAP_PRESENT_TOFROM =          (GOMP_MAP_LAST | 7)
   };
 
 #define GOMP_MAP_COPY_TO_P(X) \
@@ -243,7 +241,8 @@ enum gomp_map_kind
   (((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_FLAG_FORCE)
 
 #define GOMP_MAP_PRESENT_P(X) \
-  (((X) & GOMP_MAP_FLAG_PRESENT) == GOMP_MAP_FLAG_PRESENT)
+  (((X) & GOMP_MAP_FLAG_PRESENT) == GOMP_MAP_FLAG_PRESENT \
+   || (X) == GOMP_MAP_FORCE_PRESENT)
 
 
 /* Asynchronous behavior.  Keep in sync with

diff --git a/libgomp/target.c b/libgomp/target.c
index a9e8005c5881cf6e8fb60326fc8f980ded5a0e11..e3c4121a09ff4afa6f0ae77178e99b244aa912cf 100644 (file)
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -358,8 +358,8 @@ gomp_to_device_kind_p (int kind)
     case GOMP_MAP_FORCE_ALLOC:
     case GOMP_MAP_FORCE_FROM:
     case GOMP_MAP_ALWAYS_FROM:
-    case GOMP_MAP_PRESENT_FROM:
     case GOMP_MAP_ALWAYS_PRESENT_FROM:
+    case GOMP_MAP_FORCE_PRESENT:
       return false;
     default:
       return true;
@@ -1699,37 +1699,29 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
                    i = j - 1;
                    break;
                  case GOMP_MAP_FORCE_PRESENT:
+                 case GOMP_MAP_ALWAYS_PRESENT_TO:
+                 case GOMP_MAP_ALWAYS_PRESENT_FROM:
+                 case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
                    {
                      /* We already looked up the memory region above and it
                         was missing.  */
                      size_t size = k->host_end - k->host_start;
                      gomp_mutex_unlock (&devicep->lock);
 #ifdef HAVE_INTTYPES_H
-                     gomp_fatal ("present clause: !acc_is_present (%p, "
-                                 "%"PRIu64" (0x%"PRIx64"))",
-                                 (void *) k->host_start,
-                                 (uint64_t) size, (uint64_t) size);
+                     gomp_fatal ("present clause: not present on the device "
+                                 "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), "
+                                 "dev: %d)", (void *) k->host_start,
+                                 (uint64_t) size, (uint64_t) size,
+                                 devicep->target_id);
 #else
-                     gomp_fatal ("present clause: !acc_is_present (%p, "
-                                 "%lu (0x%lx))", (void *) k->host_start,
-                                 (unsigned long) size, (unsigned long) size);
+                     gomp_fatal ("present clause: not present on the device "
+                                 "(addr: %p, size: %lu (0x%lx), dev: %d)",
+                                 (void *) k->host_start,
+                                 (unsigned long) size, (unsigned long) size,
+                                 devicep->target_id);
 #endif
                    }
                    break;
-                 case GOMP_MAP_PRESENT_ALLOC:
-                 case GOMP_MAP_PRESENT_TO:
-                 case GOMP_MAP_PRESENT_FROM:
-                 case GOMP_MAP_PRESENT_TOFROM:
-                 case GOMP_MAP_ALWAYS_PRESENT_TO:
-                 case GOMP_MAP_ALWAYS_PRESENT_FROM:
-                 case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
-                   /* We already looked up the memory region above and it
-                      was missing.  */
-                   gomp_mutex_unlock (&devicep->lock);
-                   gomp_fatal ("present clause: not present on the device "
-                               "(%p, %d)",
-                               (void *) k->host_start, devicep->target_id);
-                   break;
                  case GOMP_MAP_FORCE_DEVICEPTR:
                    assert (k->host_end - k->host_start == sizeof (void *));
                    gomp_copy_host2dev (devicep, aq,
@@ -2149,9 +2141,18 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
                /* We already looked up the memory region above and it
                   was missing.  */
                gomp_mutex_unlock (&devicep->lock);
+#ifdef HAVE_INTTYPES_H
                gomp_fatal ("present clause: not present on the device "
-                           "(%p, %d)",
-                           (void *) hostaddrs[i], devicep->target_id);
+                           "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), "
+                           "dev: %d)", (void *) hostaddrs[i],
+                           (uint64_t) sizes[i], (uint64_t) sizes[i],
+                           devicep->target_id);
+#else
+               gomp_fatal ("present clause: not present on the device "
+                           "(addr: %p, size: %lu (0x%lx), dev: %d)",
+                           (void *) hostaddrs[i], (unsigned long) sizes[i],
+                           (unsigned long) sizes[i], devicep->target_id);
+#endif
              }
          }
       }
@@ -3465,9 +3466,7 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
              case GOMP_MAP_FORCE_TOFROM:
              case GOMP_MAP_ALWAYS_TO:
              case GOMP_MAP_ALWAYS_TOFROM:
-             case GOMP_MAP_PRESENT_FROM:
-             case GOMP_MAP_PRESENT_TO:
-             case GOMP_MAP_PRESENT_TOFROM:
+             case GOMP_MAP_FORCE_PRESENT:
              case GOMP_MAP_ALWAYS_PRESENT_FROM:
              case GOMP_MAP_ALWAYS_PRESENT_TO:
              case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
@@ -3710,8 +3709,6 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
              case GOMP_MAP_FORCE_TOFROM:
              case GOMP_MAP_ALWAYS_FROM:
              case GOMP_MAP_ALWAYS_TOFROM:
-             case GOMP_MAP_PRESENT_FROM:
-             case GOMP_MAP_PRESENT_TOFROM:
              case GOMP_MAP_ALWAYS_PRESENT_FROM:
              case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
                copy = true;

diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
index aa343197e352c0bb473145db37519e5cb1975ab6..5eaa9cd1ad31865c8460bdeb62dcc8780215cd09 100644 (file)
--- a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
@@ -4,24 +4,34 @@
 
 int main (void)
 {
-  int a[N], b[N], c[N];
+  int a[N], b[N], c[N], d[N];
 
   for (int i = 0; i < N; i++) {
     a[i] = i * 2;
     b[i] = i * 3 + 1;
+    d[i] = i * 5;
   }
 
-  #pragma omp target enter data map (alloc: a, c)
-    /* a has already been allocated, so this should be okay.  */
-    #pragma omp target map (present, to: a)
+  #pragma omp target enter data map (alloc: c, d) map(to: a)
+    #pragma omp target map (present, always, to: d)
+      for (int i = 0; i < N; i++)
+       if (d[i] != i * 5)
+         __builtin_abort ();
+
+    /* a has already been mapped and 'c' allocated so this should be okay.  */
+    #pragma omp target map (present, to: a) map(present, always, from: c)
       for (int i = 0; i < N; i++)
        c[i] = a[i];
 
+    for (int i = 0; i < N; i++)
+      if (c[i] != i * 2)
+       __builtin_abort ();
+
     fprintf (stderr, "CheCKpOInT\n");
     /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
 
     /* b has not been allocated, so this should result in an error.  */
-    /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
+    /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
     /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */
     #pragma omp target map (present, to: b)
       for (int i = 0; i < N; i++)

diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
index ad11023b2d677833481ecfd9b2b85d6ad51e711e..07ae90b5aaaf1402517370d5520c2e37c7113594 100644 (file)
--- a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
@@ -21,7 +21,7 @@ int main (void)
     /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
 
     /* b has not been allocated, so this should result in an error.  */
-    /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
+    /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
     /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */
     #pragma omp target defaultmap (present)
       for (int i = 0; i < N; i++)

diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
index 455519af4058620f13dc777b7d231c41ee953cfb..582247dc05e15e43863a684670866e37c63b064c 100644 (file)
--- a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
@@ -16,11 +16,24 @@ int main (void)
   /* This should work as a has already been allocated.  */
   #pragma omp target update to (present: a)
 
+  #pragma omp target map(present,alloc: a, c)
+  for (int i = 0; i < N; i++) {
+    if (a[i] != i * 2)
+      __builtin_abort ();
+    c[i] = 23*i;
+  }
+
+  #pragma omp target update from(present : c)
+  for (int i = 0; i < N; i++) {
+    if (c[i] != 23*i)
+      __builtin_abort ();
+  }
+
   fprintf (stderr, "CheCKpOInT\n");
   /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
 
   /* This should fail as b has not been allocated.  */
-  /* { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
+    /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */
   /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */
   #pragma omp target update to (present: b)
 

diff --git a/libgomp/testsuite/libgomp.fortran/target-present-1.f90 b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
index 768166fcff7259e62b08ac664ef813f75b0c92e7..fc13609d52835efffcd49605875ee4e92d8b3a0a 100644 (file)
--- a/libgomp/testsuite/libgomp.fortran/target-present-1.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
@@ -20,7 +20,7 @@ program main
     ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
 
     ! b has not been allocated, so this should result in an error.
-    ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } }
+    ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } }
     ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } }
     !$omp target map (present, to: b)
       do i = 1, N

diff --git a/libgomp/testsuite/libgomp.fortran/target-present-2.f90 b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
index 8f2c24ef5f261ad369c488c107b30844f1f77b57..524d01d94656219a36f5cc310a16f818d2cfe6f6 100644 (file)
--- a/libgomp/testsuite/libgomp.fortran/target-present-2.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
@@ -20,7 +20,7 @@ program main
     ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
 
     ! b has not been allocated, so this should result in an error.
-    ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } }
+    ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } }
     ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } }
     !$omp target defaultmap (present)
       do i = 1, N

diff --git a/libgomp/testsuite/libgomp.fortran/target-present-3.f90 b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
index eb29c907624677dd9e620075493514937b47e403..dd4af4c16511ae16aa2eeddaf29b39f74ee914c2 100644 (file)
--- a/libgomp/testsuite/libgomp.fortran/target-present-3.f90
+++ b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
@@ -9,14 +9,27 @@ program main
   end do
 
   !$omp target enter data map (alloc: a, c)
-    ! This should work as a has already been allocated.
-    !$omp target update to (present: a)
+
+  ! This should work as a has already been allocated.
+  !$omp target update to (present: a)
+
+  !$omp target map(present, alloc: a, c)
+    do i = 1, N
+      if (a(i) /= i * 2) stop 1
+      c(i) = 23 * i
+    end do
+  !$omp end target
+
+  !$omp target update from (present: c)
+    do i = 1, N
+      if (c(i) /= 23 * i) stop 1
+    end do
 
     print *, "CheCKpOInT"
     ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
 
     ! This should fail as b has not been allocated.
-    ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" { target offload_device_nonshared_as } }
+    ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } }
     ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } }
     !$omp target update to (present: b)
   !$omp target exit data map (from: c)

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c
index 61c8109a7e0ca8fd7dcc34910e7f7ca5d131f7c6..02fbfdaf9ed10a912fb47e95931a9a7e4ae94593 100644 (file)
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c
@@ -48,5 +48,5 @@ main (int argc, char **argv)
 }
 
 /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "present clause: !acc_is_present" } */
+/* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" } */
 /* { dg-shouldfail "" } */