public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-13] openmp: Add support for the 'present' modifier
@ 2023-06-09 11:18 Tobias Burnus
0 siblings, 0 replies; 2+ messages in thread
From: Tobias Burnus @ 2023-06-09 11:18 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:167b4964156162bfe103d088cf38973f495f54bd
commit 167b4964156162bfe103d088cf38973f495f54bd
Author: Kwok Cheung Yeung <kcy@codesourcery.com>
Date: Fri Jun 9 12:43:55 2023 +0200
openmp: Add support for the 'present' modifier
This implements support for the OpenMP 5.1 'present' modifier, which can be
used in map clauses in the 'target', 'target data', 'target data enter' and
'target data exit' constructs, and in the 'to' and 'from' clauses of the
'target update' construct. It is also supported in defaultmap.
The modifier triggers a fatal runtime error if the data specified by the
clause is not already present on the target device. It can also be combined
with 'always' in map clauses.
2023-06-06 Kwok Cheung Yeung <kcy@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
gcc/c/
* c-parser.cc (c_parser_omp_clause_defaultmap,
c_parser_omp_clause_map): Parse 'present'.
(c_parser_omp_clause_to, c_parser_omp_clause_from): Remove.
(c_parser_omp_clause_from_to): New; parse to/from clauses with
optional present modifer.
(c_parser_omp_all_clauses): Update call.
(c_parser_omp_target_data, c_parser_omp_target_enter_data,
c_parser_omp_target_exit_data): Handle new map enum values
for 'present' mapping.
gcc/cp/
* parser.cc (cp_parser_omp_clause_defaultmap,
cp_parser_omp_clause_map): Parse 'present'.
(cp_parser_omp_clause_from_to): New; parse to/from
clauses with optional 'present' modifier.
(cp_parser_omp_all_clauses): Update call.
(cp_parser_omp_target_data, cp_parser_omp_target_enter_data,
cp_parser_omp_target_exit_data): Handle new enum value for
'present' mapping.
* semantics.cc (finish_omp_target): Likewise.
gcc/fortran/
* dump-parse-tree.cc (show_omp_namelist): Display 'present' map
modifier.
(show_omp_clauses): Display 'present' motion modifier for 'to'
and 'from' clauses.
* gfortran.h (enum gfc_omp_map_op): Add entries with 'present'
modifiers.
(struct gfc_omp_namelist): Add 'present_modifer'.
* openmp.cc (gfc_match_motion_var_list): New, handles optional
'present' modifier for to/from clauses.
(gfc_match_omp_clauses): Call it for to/from clauses; parse 'present'
in defaultmap and map clauses.
(resolve_omp_clauses): Allow 'present' modifiers on 'target',
'target data', 'target enter' and 'target exit' directives.
* trans-openmp.cc (gfc_trans_omp_clauses): Apply 'present' modifiers
to tree node for 'map', 'to' and 'from' clauses. Apply 'present' for
defaultmap.
gcc/
* gimplify.cc (omp_notice_variable): Apply GOVD_MAP_ALLOC_ONLY flag
and defaultmap flags if the defaultmap has GOVD_MAP_FORCE_PRESENT flag
set.
(omp_get_attachment): Handle map clauses with 'present' modifier.
(omp_group_base): Likewise.
(gimplify_scan_omp_clauses): Reorder present maps to come first.
Set GOVD flags for present defaultmaps.
(gimplify_adjust_omp_clauses_1): Set map kind for present defaultmaps.
* omp-low.cc (scan_sharing_clauses): Handle 'always, present' map
clauses.
(lower_omp_target): Handle map clauses with 'present' modifier.
Handle 'to' and 'from' clauses with 'present'.
* tree-core.h (enum omp_clause_defaultmap_kind): Add
OMP_CLAUSE_DEFAULTMAP_PRESENT defaultmap kind.
* tree-pretty-print.cc (dump_omp_clause): Handle 'map', 'to' and
'from' clauses with 'present' modifier. Handle present defaultmap.
* tree.h (OMP_CLAUSE_MOTION_PRESENT): New #define.
include/
* gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_5): New.
(GOMP_MAP_FLAG_FORCE): Redefine.
(GOMP_MAP_FLAG_PRESENT, GOMP_MAP_FLAG_ALWAYS_PRESENT): New.
(enum gomp_map_kind): Add map kinds with 'present' modifiers.
(GOMP_MAP_COPY_TO_P, GOMP_MAP_COPY_FROM_P): Evaluate to true for
map variants with 'present'
(GOMP_MAP_ALWAYS_TO_P, GOMP_MAP_ALWAYS_FROM_P): Evaluate to true
for map variants with 'always, present' modifiers.
(GOMP_MAP_ALWAYS): Redefine.
(GOMP_MAP_FORCE_P, GOMP_MAP_PRESENT_P): New.
libgomp/
* libgomp.texi (OpenMP 5.1 Impl. status): Set 'present' support for
defaultmap to 'Y', add 'Y' entry for 'present' on to/from/map clauses.
* target.c (gomp_to_device_kind_p): Add map kinds with 'present'
modifier.
(gomp_map_vars_existing): Use new GOMP_MAP_FORCE_P macro.
(gomp_map_vars_internal, gomp_update, gomp_target_rev):
Emit runtime error if memory region not present.
* testsuite/libgomp.c-c++-common/target-present-1.c: New test.
* testsuite/libgomp.c-c++-common/target-present-2.c: New test.
* testsuite/libgomp.c-c++-common/target-present-3.c: New test.
* testsuite/libgomp.fortran/target-present-1.f90: New test.
* testsuite/libgomp.fortran/target-present-2.f90: New test.
* testsuite/libgomp.fortran/target-present-3.f90: New test.
gcc/testsuite/
* c-c++-common/gomp/map-6.c: Update dg-error, extend to test for
duplicated 'present' and extend scan-dump tests for 'present'.
* gfortran.dg/gomp/defaultmap-1.f90: Update dg-error.
* gfortran.dg/gomp/map-7.f90: Extend parse and dump test for
'present'.
* gfortran.dg/gomp/map-8.f90: Extend for duplicate 'present'
modifier checking.
* c-c++-common/gomp/defaultmap-4.c: New test.
* c-c++-common/gomp/map-9.c: New test.
* c-c++-common/gomp/target-update-1.c: New test.
* gfortran.dg/gomp/defaultmap-8.f90: New test.
* gfortran.dg/gomp/map-11.f90: New test.
* gfortran.dg/gomp/map-12.f90: New test.
* gfortran.dg/gomp/target-update-1.f90: New test.
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>
(cherry picked from commit 4ede915d5dde935a16df2c6640aee5ab22348d30)
Diff:
---
gcc/ChangeLog.omp | 24 ++++
gcc/c/ChangeLog.omp | 16 +++
gcc/c/c-parser.cc | 126 +++++++++++++++++----
gcc/cp/ChangeLog.omp | 16 +++
gcc/cp/parser.cc | 115 +++++++++++++++++--
gcc/cp/semantics.cc | 7 ++
gcc/fortran/ChangeLog.omp | 24 ++++
gcc/fortran/dump-parse-tree.cc | 14 +++
gcc/fortran/gfortran.h | 8 ++
gcc/fortran/openmp.cc | 98 +++++++++++++---
gcc/fortran/trans-openmp.cc | 42 ++++++-
gcc/gimplify.cc | 69 +++++++++++
gcc/omp-low.cc | 26 ++++-
gcc/testsuite/ChangeLog.omp | 22 ++++
gcc/testsuite/c-c++-common/gomp/clauses-2.c | 2 +-
gcc/testsuite/c-c++-common/gomp/defaultmap-4.c | 24 ++++
gcc/testsuite/c-c++-common/gomp/map-6.c | 70 +++++++++++-
gcc/testsuite/c-c++-common/gomp/map-9.c | 32 ++++++
gcc/testsuite/c-c++-common/gomp/target-update-1.c | 15 +++
gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90 | 2 +-
gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 | 26 +++++
gcc/testsuite/gfortran.dg/gomp/map-11.f90 | 34 ++++++
gcc/testsuite/gfortran.dg/gomp/map-12.f90 | 67 +++++++++++
gcc/testsuite/gfortran.dg/gomp/map-7.f90 | 30 ++++-
gcc/testsuite/gfortran.dg/gomp/map-8.f90 | 11 ++
gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 | 13 +++
gcc/tree-core.h | 3 +-
gcc/tree-pretty-print.cc | 28 +++++
gcc/tree.h | 3 +
include/ChangeLog.omp | 17 +++
include/gomp-constants.h | 47 ++++++--
libgomp/ChangeLog.omp | 20 ++++
libgomp/libgomp.texi | 4 +-
libgomp/target.c | 66 ++++++++++-
.../libgomp.c-c++-common/target-present-1.c | 27 +++++
.../libgomp.c-c++-common/target-present-2.c | 27 +++++
.../libgomp.c-c++-common/target-present-3.c | 27 +++++
.../testsuite/libgomp.fortran/target-present-1.f90 | 30 +++++
.../testsuite/libgomp.fortran/target-present-2.f90 | 30 +++++
.../testsuite/libgomp.fortran/target-present-3.f90 | 22 ++++
40 files changed, 1213 insertions(+), 71 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index d4c4e5fe3c5..e59213c3e7d 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,27 @@
+2023-06-07 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from mainline:
+ 2023-06-06 Kwok Cheung Yeung <kcy@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+
+ * gimplify.cc (omp_notice_variable): Apply GOVD_MAP_ALLOC_ONLY flag
+ and defaultmap flags if the defaultmap has GOVD_MAP_FORCE_PRESENT flag
+ set.
+ (omp_get_attachment): Handle map clauses with 'present' modifier.
+ (omp_group_base): Likewise.
+ (gimplify_scan_omp_clauses): Reorder present maps to come first.
+ Set GOVD flags for present defaultmaps.
+ (gimplify_adjust_omp_clauses_1): Set map kind for present defaultmaps.
+ * omp-low.cc (scan_sharing_clauses): Handle 'always, present' map
+ clauses.
+ (lower_omp_target): Handle map clauses with 'present' modifier.
+ Handle 'to' and 'from' clauses with 'present'.
+ * tree-core.h (enum omp_clause_defaultmap_kind): Add
+ OMP_CLAUSE_DEFAULTMAP_PRESENT defaultmap kind.
+ * tree-pretty-print.cc (dump_omp_clause): Handle 'map', 'to' and
+ 'from' clauses with 'present' modifier. Handle present defaultmap.
+ * tree.h (OMP_CLAUSE_MOTION_PRESENT): New #define.
+
2023-06-07 Tobias Burnus <tobias@codesourcery.com>
Reverted:
diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index 960f56399d7..f2111ef1e5e 100644
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,19 @@
+2023-06-07 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from mainline:
+ 2023-06-06 Kwok Cheung Yeung <kcy@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+
+ * c-parser.cc (c_parser_omp_clause_defaultmap,
+ c_parser_omp_clause_map): Parse 'present'.
+ (c_parser_omp_clause_to, c_parser_omp_clause_from): Remove.
+ (c_parser_omp_clause_from_to): New; parse to/from clauses with
+ optional present modifer.
+ (c_parser_omp_all_clauses): Update call.
+ (c_parser_omp_target_data, c_parser_omp_target_enter_data,
+ c_parser_omp_target_exit_data): Handle new map enum values
+ for 'present' mapping.
+
2023-06-07 Tobias Burnus <tobias@codesourcery.com>
Reverted:
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 6e3401b6532..eb56a9afc4e 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -14970,6 +14970,13 @@ c_parser_omp_clause_defaultmap (c_parser *parser, tree list)
goto invalid_behavior;
break;
+ case 'p':
+ if (strcmp ("present", p) == 0)
+ behavior = OMP_CLAUSE_DEFAULTMAP_PRESENT;
+ else
+ goto invalid_behavior;
+ break;
+
case 't':
if (strcmp ("tofrom", p) == 0)
behavior = OMP_CLAUSE_DEFAULTMAP_TOFROM;
@@ -17348,6 +17355,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
int always_modifier = 0;
int close_modifier = 0;
+ int present_modifier = 0;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
c_token *tok = c_parser_peek_token (parser);
@@ -17379,11 +17387,21 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
}
close_modifier++;
}
+ else if (strcmp ("present", p) == 0)
+ {
+ if (present_modifier)
+ {
+ c_parser_error (parser, "too many %<present%> modifiers");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ present_modifier++;
+ }
else
{
c_parser_error (parser, "%<#pragma omp target%> with "
- "modifier other than %<always%> or "
- "%<close%> on %<map%> clause");
+ "modifier other than %<always%>, %<close%> "
+ "or %<present%> on %<map%> clause");
parens.skip_until_found_close (parser);
return list;
}
@@ -17395,14 +17413,25 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
&& c_parser_peek_2nd_token (parser)->type == CPP_COLON)
{
const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ int always_present_modifier = always_modifier && present_modifier;
+
if (strcmp ("alloc", p) == 0)
- kind = GOMP_MAP_ALLOC;
+ kind = present_modifier ? GOMP_MAP_PRESENT_ALLOC : GOMP_MAP_ALLOC;
else if (strcmp ("to", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+ kind = (always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TO
+ : present_modifier ? GOMP_MAP_PRESENT_TO
+ : always_modifier ? GOMP_MAP_ALWAYS_TO
+ : GOMP_MAP_TO);
else if (strcmp ("from", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+ kind = (always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_FROM
+ : present_modifier ? GOMP_MAP_PRESENT_FROM
+ : always_modifier ? GOMP_MAP_ALWAYS_FROM
+ : GOMP_MAP_FROM);
else if (strcmp ("tofrom", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+ kind = (always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TOFROM
+ : present_modifier ? GOMP_MAP_PRESENT_TOFROM
+ : always_modifier ? GOMP_MAP_ALWAYS_TOFROM
+ : GOMP_MAP_TOFROM);
else if (strcmp ("release", p) == 0)
kind = GOMP_MAP_RELEASE;
else if (strcmp ("delete", p) == 0)
@@ -17658,23 +17687,43 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list)
}
/* OpenMP 4.0:
- to ( variable-list ) */
+ from ( variable-list )
+ to ( variable-list )
+
+ OpenMP 5.1:
+ from ( [present :] variable-list )
+ to ( [present :] variable-list ) */
static tree
-c_parser_omp_clause_to (c_parser *parser, tree list)
+c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
+ tree list)
{
- return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list,
- C_ORT_OMP, true);
-}
+ location_t loc = c_parser_peek_token (parser)->location;
+ matching_parens parens;
+ if (!parens.require_open (parser))
+ return list;
-/* OpenMP 4.0:
- from ( variable-list ) */
+ bool present = false;
+ c_token *token = c_parser_peek_token (parser);
-static tree
-c_parser_omp_clause_from (c_parser *parser, tree list)
-{
- return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list,
- C_ORT_OMP, true);
+ if (token->type == CPP_NAME
+ && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0
+ && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+ {
+ present = true;
+ c_parser_consume_token (parser);
+ c_parser_consume_token (parser);
+ }
+
+ tree nl = c_parser_omp_variable_list (parser, loc, kind, list, C_ORT_OMP,
+ true);
+ parens.skip_until_found_close (parser);
+
+ if (present)
+ for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ OMP_CLAUSE_MOTION_PRESENT (c) = 1;
+
+ return nl;
}
/* OpenMP 4.0:
@@ -18242,11 +18291,13 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = nl;
}
else
- clauses = c_parser_omp_clause_to (parser, clauses);
+ clauses = c_parser_omp_clause_from_to (parser, OMP_CLAUSE_TO,
+ clauses);
c_name = "to";
break;
case PRAGMA_OMP_CLAUSE_FROM:
- clauses = c_parser_omp_clause_from (parser, clauses);
+ clauses = c_parser_omp_clause_from_to (parser, OMP_CLAUSE_FROM,
+ clauses);
c_name = "from";
break;
case PRAGMA_OMP_CLAUSE_UNIFORM:
@@ -22125,11 +22176,18 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
@@ -22275,7 +22333,10 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
map_seen = 3;
break;
case GOMP_MAP_TOFROM:
@@ -22286,6 +22347,14 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_TO);
map_seen = 3;
break;
+ case GOMP_MAP_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_TO);
+ map_seen = 3;
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_TO);
+ map_seen = 3;
+ break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
@@ -22373,6 +22442,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
{
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
map_seen = 3;
@@ -22385,6 +22456,14 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_FROM);
map_seen = 3;
break;
+ case GOMP_MAP_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_FROM);
+ map_seen = 3;
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_FROM);
+ map_seen = 3;
+ break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
@@ -22631,11 +22710,18 @@ check_clauses:
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index 56e40aee876..1551eae13b3 100644
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,19 @@
+2023-06-07 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from mainline:
+ 2023-06-06 Kwok Cheung Yeung <kcy@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+
+ * parser.cc (cp_parser_omp_clause_defaultmap,
+ cp_parser_omp_clause_map): Parse 'present'.
+ (cp_parser_omp_clause_from_to): New; parse to/from
+ clauses with optional 'present' modifier.
+ (cp_parser_omp_all_clauses): Update call.
+ (cp_parser_omp_target_data, cp_parser_omp_target_enter_data,
+ cp_parser_omp_target_exit_data): Handle new enum value for
+ 'present' mapping.
+ * semantics.cc (finish_omp_target): Likewise.
+
2023-06-07 Tobias Burnus <tobias@codesourcery.com>
Reverted:
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index f25ae875ee9..ead672a4612 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -38764,6 +38764,13 @@ cp_parser_omp_clause_defaultmap (cp_parser *parser, tree list,
goto invalid_behavior;
break;
+ case 'p':
+ if (strcmp ("present", p) == 0)
+ behavior = OMP_CLAUSE_DEFAULTMAP_PRESENT;
+ else
+ goto invalid_behavior;
+ break;
+
case 't':
if (strcmp ("tofrom", p) == 0)
behavior = OMP_CLAUSE_DEFAULTMAP_TOFROM;
@@ -40750,6 +40757,42 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc)
return nlist;
}
+/* OpenMP 4.0:
+ from ( variable-list )
+ to ( variable-list )
+
+ OpenMP 5.1:
+ from ( [present :] variable-list )
+ to ( [present :] variable-list ) */
+
+static tree
+cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
+ tree list)
+{
+ if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
+ return list;
+
+ bool present = false;
+ cp_token *token = cp_lexer_peek_token (parser->lexer);
+
+ if (token->type == CPP_NAME
+ && strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+ {
+ present = true;
+ cp_lexer_consume_token (parser->lexer);
+ cp_lexer_consume_token (parser->lexer);
+ }
+
+ tree nl = cp_parser_omp_var_list_no_open (parser, kind, list, NULL, C_ORT_OMP,
+ true);
+ if (present)
+ for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ OMP_CLAUSE_MOTION_PRESENT (c) = 1;
+
+ return nl;
+}
+
/* OpenMP 4.0:
map ( map-kind : variable-list )
map ( variable-list )
@@ -40796,6 +40839,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
bool always_modifier = false;
bool close_modifier = false;
+ bool present_modifier = false;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
cp_token *tok = cp_lexer_peek_token (parser->lexer);
@@ -40832,11 +40876,24 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
}
close_modifier = true;
}
+ else if (strcmp ("present", p) == 0)
+ {
+ if (present_modifier)
+ {
+ cp_parser_error (parser, "too many %<present%> modifiers");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ present_modifier = true;
+ }
else
{
cp_parser_error (parser, "%<#pragma omp target%> with "
- "modifier other than %<always%> or "
- "%<close%> on %<map%> clause");
+ "modifier other than %<always%>, %<close%> "
+ "or %<present%> on %<map%> clause");
cp_parser_skip_to_closing_parenthesis (parser,
/*recovering=*/true,
/*or_comma=*/false,
@@ -40852,15 +40909,25 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
{
tree id = cp_lexer_peek_token (parser->lexer)->u.value;
const char *p = IDENTIFIER_POINTER (id);
+ int always_present_modifier = always_modifier && present_modifier;
if (strcmp ("alloc", p) == 0)
- kind = GOMP_MAP_ALLOC;
+ kind = present_modifier ? GOMP_MAP_PRESENT_ALLOC : GOMP_MAP_ALLOC;
else if (strcmp ("to", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+ kind = (always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TO
+ : present_modifier ? GOMP_MAP_PRESENT_TO
+ : always_modifier ? GOMP_MAP_ALWAYS_TO
+ : GOMP_MAP_TO);
else if (strcmp ("from", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+ kind = (always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_FROM
+ : present_modifier ? GOMP_MAP_PRESENT_FROM
+ : always_modifier ? GOMP_MAP_ALWAYS_FROM
+ : GOMP_MAP_FROM);
else if (strcmp ("tofrom", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+ kind = (always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TOFROM
+ : present_modifier ? GOMP_MAP_PRESENT_TOFROM
+ : always_modifier ? GOMP_MAP_ALWAYS_TOFROM
+ : GOMP_MAP_TOFROM);
else if (strcmp ("release", p) == 0)
kind = GOMP_MAP_RELEASE;
else
@@ -41637,13 +41704,13 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = nl;
}
else
- clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses,
- C_ORT_OMP, true);
+ clauses = cp_parser_omp_clause_from_to (parser, OMP_CLAUSE_TO,
+ clauses);
c_name = "to";
break;
case PRAGMA_OMP_CLAUSE_FROM:
- clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses,
- C_ORT_OMP, true);
+ clauses = cp_parser_omp_clause_from_to (parser, OMP_CLAUSE_FROM,
+ clauses);
c_name = "from";
break;
case PRAGMA_OMP_CLAUSE_UNIFORM:
@@ -45572,11 +45639,18 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
@@ -45679,7 +45753,10 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
map_seen = 3;
break;
case GOMP_MAP_TOFROM:
@@ -45690,6 +45767,14 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_TO);
map_seen = 3;
break;
+ case GOMP_MAP_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_TO);
+ map_seen = 3;
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_TO);
+ map_seen = 3;
+ break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
@@ -45782,6 +45867,8 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
{
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
map_seen = 3;
@@ -45794,6 +45881,14 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_FROM);
map_seen = 3;
break;
+ case GOMP_MAP_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_FROM);
+ map_seen = 3;
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_FROM);
+ map_seen = 3;
+ break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index f86e8281d3d..f5c9521c653 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -10282,11 +10282,18 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index d94c6596b42..28fed83a33a 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,27 @@
+2023-06-07 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from mainline:
+ 2023-06-06 Kwok Cheung Yeung <kcy@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+
+ * dump-parse-tree.cc (show_omp_namelist): Display 'present' map
+ modifier.
+ (show_omp_clauses): Display 'present' motion modifier for 'to'
+ and 'from' clauses.
+
+ * gfortran.h (enum gfc_omp_map_op): Add entries with 'present'
+ modifiers.
+ (struct gfc_omp_namelist): Add 'present_modifer'.
+ * openmp.cc (gfc_match_motion_var_list): New, handles optional
+ 'present' modifier for to/from clauses.
+ (gfc_match_omp_clauses): Call it for to/from clauses; parse 'present'
+ in defaultmap and map clauses.
+ (resolve_omp_clauses): Allow 'present' modifiers on 'target',
+ 'target data', 'target enter' and 'target exit' directives.
+ * trans-openmp.cc (gfc_trans_omp_clauses): Apply 'present' modifiers
+ to tree node for 'map', 'to' and 'from' clauses. Apply 'present' for
+ defaultmap.
+
2023-06-07 Tobias Burnus <tobias@codesourcery.com>
Reverted:
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index 266a4c7b69f..ae1673a3de6 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -1453,9 +1453,20 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
case OMP_MAP_TO: fputs ("to:", dumpfile); break;
case OMP_MAP_FROM: fputs ("from:", dumpfile); break;
case OMP_MAP_TOFROM: fputs ("tofrom:", dumpfile); break;
+ case OMP_MAP_PRESENT_ALLOC: fputs ("present,alloc:", dumpfile); break;
+ case OMP_MAP_PRESENT_TO: fputs ("present,to:", dumpfile); break;
+ case OMP_MAP_PRESENT_FROM: fputs ("present,from:", dumpfile); break;
+ case OMP_MAP_PRESENT_TOFROM:
+ fputs ("present,tofrom:", dumpfile); break;
case OMP_MAP_ALWAYS_TO: fputs ("always,to:", dumpfile); break;
case OMP_MAP_ALWAYS_FROM: fputs ("always,from:", dumpfile); break;
case OMP_MAP_ALWAYS_TOFROM: fputs ("always,tofrom:", dumpfile); break;
+ case OMP_MAP_ALWAYS_PRESENT_TO:
+ fputs ("always,present,to:", dumpfile); break;
+ case OMP_MAP_ALWAYS_PRESENT_FROM:
+ fputs ("always,present,from:", dumpfile); break;
+ case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+ fputs ("always,present,tofrom:", dumpfile); break;
case OMP_MAP_DELETE: fputs ("delete:", dumpfile); break;
case OMP_MAP_RELEASE: fputs ("release:", dumpfile); break;
default: break;
@@ -1793,6 +1804,9 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
fputs ("inscan, ", dumpfile);
if (list_type == OMP_LIST_REDUCTION_TASK)
fputs ("task, ", dumpfile);
+ if ((list_type == OMP_LIST_TO || list_type == OMP_LIST_FROM)
+ && omp_clauses->lists[list_type]->u.present_modifier)
+ fputs ("present:", dumpfile);
show_omp_namelist (list_type, omp_clauses->lists[list_type]);
fputc (')', dumpfile);
}
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index e2089c29892..9c8d88669a0 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1310,6 +1310,13 @@ enum gfc_omp_map_op
OMP_MAP_ALWAYS_TO,
OMP_MAP_ALWAYS_FROM,
OMP_MAP_ALWAYS_TOFROM,
+ OMP_MAP_PRESENT_ALLOC,
+ OMP_MAP_PRESENT_TO,
+ OMP_MAP_PRESENT_FROM,
+ OMP_MAP_PRESENT_TOFROM,
+ OMP_MAP_ALWAYS_PRESENT_TO,
+ OMP_MAP_ALWAYS_PRESENT_FROM,
+ OMP_MAP_ALWAYS_PRESENT_TOFROM,
OMP_MAP_DECLARE_ALLOCATE,
OMP_MAP_DECLARE_DEALLOCATE
};
@@ -1365,6 +1372,7 @@ typedef struct gfc_omp_namelist
} linear;
struct gfc_common_head *common;
bool lastprivate_conditional;
+ bool present_modifier;
} u;
union
{
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 413ee2cddba..ca9a8e665d1 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -1426,6 +1426,30 @@ failed:
return MATCH_NO;
}
+/* Match target update's to/from( [present:] var-list). */
+
+static match
+gfc_match_motion_var_list (const char *str, gfc_omp_namelist **list,
+ gfc_omp_namelist ***headp)
+{
+ match m = gfc_match (str);
+ if (m != MATCH_YES)
+ return m;
+
+ match m_present = gfc_match (" present : ");
+
+ m = gfc_match_omp_variable_list ("", list, false, NULL, headp, true, true);
+ if (m != MATCH_YES)
+ return m;
+ if (m_present == MATCH_YES)
+ {
+ gfc_omp_namelist *n;
+ for (n = **headp; n; n = n->next)
+ n->u.present_modifier = true;
+ }
+ return MATCH_YES;
+}
+
/* reduction ( reduction-modifier, reduction-operator : variable-list )
in_reduction ( reduction-operator : variable-list )
task_reduction ( reduction-operator : variable-list ) */
@@ -2708,6 +2732,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
behavior = OMP_DEFAULTMAP_FROM;
else if (gfc_match ("firstprivate ") == MATCH_YES)
behavior = OMP_DEFAULTMAP_FIRSTPRIVATE;
+ else if (gfc_match ("present ") == MATCH_YES)
+ behavior = OMP_DEFAULTMAP_PRESENT;
else if (gfc_match ("none ") == MATCH_YES)
behavior = OMP_DEFAULTMAP_NONE;
else if (gfc_match ("default ") == MATCH_YES)
@@ -2715,7 +2741,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
else
{
gfc_error ("Expected ALLOC, TO, FROM, TOFROM, FIRSTPRIVATE, "
- "NONE or DEFAULT at %C");
+ "PRESENT, NONE or DEFAULT at %C");
break;
}
if (')' == gfc_peek_ascii_char ())
@@ -3139,10 +3165,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
true) == MATCH_YES)
continue;
if ((mask & OMP_CLAUSE_FROM)
- && (gfc_match_omp_variable_list ("from (",
- &c->lists[OMP_LIST_FROM], false,
- NULL, &head, true, true)
- == MATCH_YES))
+ && gfc_match_motion_var_list ("from (", &c->lists[OMP_LIST_FROM],
+ &head) == MATCH_YES)
continue;
if ((mask & OMP_CLAUSE_UNROLL_FULL)
&& (m = gfc_match_dupl_check (!c->unroll_full, "full"))
@@ -3507,8 +3531,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
locus old_loc2 = gfc_current_locus;
int always_modifier = 0;
int close_modifier = 0;
+ int present_modifier = 0;
locus second_always_locus = old_loc2;
locus second_close_locus = old_loc2;
+ locus second_present_locus = old_loc2;
for (;;)
{
@@ -3523,20 +3549,38 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if (close_modifier++ == 1)
second_close_locus = current_locus;
}
+ else if (gfc_match ("present ") == MATCH_YES)
+ {
+ if (present_modifier++ == 1)
+ second_present_locus = current_locus;
+ }
else
break;
gfc_match (", ");
}
gfc_omp_map_op map_op = OMP_MAP_TOFROM;
+ int always_present_modifier
+ = always_modifier && present_modifier;
+
if (gfc_match ("alloc : ") == MATCH_YES)
- map_op = OMP_MAP_ALLOC;
+ map_op = (present_modifier ? OMP_MAP_PRESENT_ALLOC
+ : OMP_MAP_ALLOC);
else if (gfc_match ("tofrom : ") == MATCH_YES)
- map_op = always_modifier ? OMP_MAP_ALWAYS_TOFROM : OMP_MAP_TOFROM;
+ map_op = (always_present_modifier ? OMP_MAP_ALWAYS_PRESENT_TOFROM
+ : present_modifier ? OMP_MAP_PRESENT_TOFROM
+ : always_modifier ? OMP_MAP_ALWAYS_TOFROM
+ : OMP_MAP_TOFROM);
else if (gfc_match ("to : ") == MATCH_YES)
- map_op = always_modifier ? OMP_MAP_ALWAYS_TO : OMP_MAP_TO;
+ map_op = (always_present_modifier ? OMP_MAP_ALWAYS_PRESENT_TO
+ : present_modifier ? OMP_MAP_PRESENT_TO
+ : always_modifier ? OMP_MAP_ALWAYS_TO
+ : OMP_MAP_TO);
else if (gfc_match ("from : ") == MATCH_YES)
- map_op = always_modifier ? OMP_MAP_ALWAYS_FROM : OMP_MAP_FROM;
+ map_op = (always_present_modifier ? OMP_MAP_ALWAYS_PRESENT_FROM
+ : present_modifier ? OMP_MAP_PRESENT_FROM
+ : always_modifier ? OMP_MAP_ALWAYS_FROM
+ : OMP_MAP_FROM);
else if (gfc_match ("release : ") == MATCH_YES)
map_op = OMP_MAP_RELEASE;
else if (gfc_match ("delete : ") == MATCH_YES)
@@ -3560,6 +3604,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
&second_close_locus);
break;
}
+ if (present_modifier > 1)
+ {
+ gfc_error ("too many %<present%> modifiers at %L",
+ &second_present_locus);
+ break;
+ }
head = NULL;
if (gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP],
@@ -4112,10 +4162,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
continue;
}
else if ((mask & OMP_CLAUSE_TO)
- && (gfc_match_omp_variable_list ("to (",
- &c->lists[OMP_LIST_TO], false,
- NULL, &head, true, true)
- == MATCH_YES))
+ && gfc_match_motion_var_list ("to (", &c->lists[OMP_LIST_TO],
+ &head) == MATCH_YES)
continue;
break;
case 'u':
@@ -8745,11 +8793,18 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
{
case OMP_MAP_TO:
case OMP_MAP_ALWAYS_TO:
+ case OMP_MAP_PRESENT_TO:
+ case OMP_MAP_ALWAYS_PRESENT_TO:
case OMP_MAP_FROM:
case OMP_MAP_ALWAYS_FROM:
+ case OMP_MAP_PRESENT_FROM:
+ case OMP_MAP_ALWAYS_PRESENT_FROM:
case OMP_MAP_TOFROM:
case OMP_MAP_ALWAYS_TOFROM:
+ case OMP_MAP_PRESENT_TOFROM:
+ case OMP_MAP_ALWAYS_PRESENT_TOFROM:
case OMP_MAP_ALLOC:
+ case OMP_MAP_PRESENT_ALLOC:
break;
default:
gfc_error ("TARGET%s with map-type other than TO, "
@@ -8765,7 +8820,10 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
{
case OMP_MAP_TO:
case OMP_MAP_ALWAYS_TO:
+ case OMP_MAP_PRESENT_TO:
+ case OMP_MAP_ALWAYS_PRESENT_TO:
case OMP_MAP_ALLOC:
+ case OMP_MAP_PRESENT_ALLOC:
break;
case OMP_MAP_TOFROM:
n->u.map_op = OMP_MAP_TO;
@@ -8773,6 +8831,12 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
case OMP_MAP_ALWAYS_TOFROM:
n->u.map_op = OMP_MAP_ALWAYS_TO;
break;
+ case OMP_MAP_PRESENT_TOFROM:
+ n->u.map_op = OMP_MAP_PRESENT_TO;
+ break;
+ case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+ n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO;
+ break;
default:
gfc_error ("TARGET ENTER DATA with map-type other "
"than TO, TOFROM or ALLOC on MAP clause "
@@ -8785,6 +8849,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
{
case OMP_MAP_FROM:
case OMP_MAP_ALWAYS_FROM:
+ case OMP_MAP_PRESENT_FROM:
+ case OMP_MAP_ALWAYS_PRESENT_FROM:
case OMP_MAP_RELEASE:
case OMP_MAP_DELETE:
break;
@@ -8794,6 +8860,12 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
case OMP_MAP_ALWAYS_TOFROM:
n->u.map_op = OMP_MAP_ALWAYS_FROM;
break;
+ case OMP_MAP_PRESENT_TOFROM:
+ n->u.map_op = OMP_MAP_PRESENT_FROM;
+ break;
+ case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+ n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM;
+ break;
default:
gfc_error ("TARGET EXIT DATA with map-type other "
"than FROM, TOFROM, RELEASE, or DELETE on "
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index cb3b3be92ef..dd9ea2db4b8 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -2934,6 +2934,14 @@ gfc_omp_deep_map_kind_p (tree clause)
switch (OMP_CLAUSE_MAP_KIND (clause))
{
+ case GOMP_MAP_PRESENT_ALLOC:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
+ return false;
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
@@ -2942,11 +2950,6 @@ gfc_omp_deep_map_kind_p (tree clause)
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_FIRSTPRIVATE:
case GOMP_MAP_ALLOC:
-<<<<<<< HEAD
- case GOMP_MAP_PRESENT_ALLOC:
- return true;
-=======
->>>>>>> parent of 6e3816fa47c (openmp: Add support for the 'present' modifier)
case GOMP_MAP_POINTER:
case GOMP_MAP_TO_PSET:
case GOMP_MAP_FORCE_PRESENT:
@@ -4580,6 +4583,30 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
always_modifier = true;
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TOFROM);
break;
+ case OMP_MAP_PRESENT_ALLOC:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_ALLOC);
+ break;
+ case OMP_MAP_PRESENT_TO:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_TO);
+ break;
+ case OMP_MAP_PRESENT_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_FROM);
+ break;
+ case OMP_MAP_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_TOFROM);
+ break;
+ case OMP_MAP_ALWAYS_PRESENT_TO:
+ always_modifier = true;
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_PRESENT_TO);
+ break;
+ case OMP_MAP_ALWAYS_PRESENT_FROM:
+ always_modifier = true;
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_PRESENT_FROM);
+ break;
+ case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+ always_modifier = true;
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_PRESENT_TOFROM);
+ break;
case OMP_MAP_RELEASE:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
break;
@@ -5480,6 +5507,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
}
+ if (n->u.present_modifier)
+ OMP_CLAUSE_MOTION_PRESENT (node) = 1;
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
}
break;
@@ -6089,6 +6118,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_DEFAULTMAP_FIRSTPRIVATE:
behavior = OMP_CLAUSE_DEFAULTMAP_FIRSTPRIVATE;
break;
+ case OMP_DEFAULTMAP_PRESENT:
+ behavior = OMP_CLAUSE_DEFAULTMAP_PRESENT;
+ break;
case OMP_DEFAULTMAP_NONE: behavior = OMP_CLAUSE_DEFAULTMAP_NONE; break;
case OMP_DEFAULTMAP_DEFAULT:
behavior = OMP_CLAUSE_DEFAULTMAP_DEFAULT;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index fe6c1ba2a30..966404dab8d 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -7936,6 +7936,11 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
else if (ctx->defaultmap[gdmk]
& (GOVD_MAP_0LEN_ARRAY | GOVD_FIRSTPRIVATE))
nflags |= ctx->defaultmap[gdmk];
+ else if (ctx->defaultmap[gdmk] & GOVD_MAP_FORCE_PRESENT)
+ {
+ gcc_assert (ctx->defaultmap[gdmk] & GOVD_MAP);
+ nflags |= ctx->defaultmap[gdmk] | GOVD_MAP_ALLOC_ONLY;
+ }
else
{
gcc_assert (ctx->defaultmap[gdmk] & GOVD_MAP);
@@ -9111,6 +9116,13 @@ omp_get_attachment (omp_mapping_group *grp)
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_PRESENT_ALLOC:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
@@ -9342,6 +9354,13 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained,
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_PRESENT_ALLOC:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_NONCONTIG_ARRAY_ALLOC:
case GOMP_MAP_NONCONTIG_ARRAY_FROM:
case GOMP_MAP_NONCONTIG_ARRAY_TO:
@@ -10834,6 +10853,50 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
delete grpmap;
delete groups;
}
+
+ /* OpenMP map clauses with 'present' need to go in front of those
+ without. */
+ tree present_map_head = NULL;
+ tree *present_map_tail_p = &present_map_head;
+ tree *first_map_clause_p = NULL;
+
+ for (tree *c_p = list_p; *c_p; )
+ {
+ tree c = *c_p;
+ tree *next_c_p = &OMP_CLAUSE_CHAIN (c);
+
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+ {
+ if (!first_map_clause_p)
+ first_map_clause_p = c_p;
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_PRESENT_ALLOC:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_PRESENT_TOFROM:
+ next_c_p = c_p;
+ *c_p = OMP_CLAUSE_CHAIN (c);
+
+ OMP_CLAUSE_CHAIN (c) = NULL;
+ *present_map_tail_p = c;
+ present_map_tail_p = &OMP_CLAUSE_CHAIN (c);
+
+ break;
+
+ default:
+ break;
+ }
+ }
+
+ c_p = next_c_p;
+ }
+ if (first_map_clause_p && present_map_head)
+ {
+ tree next = *first_map_clause_p;
+ *first_map_clause_p = present_map_head;
+ *present_map_tail_p = next;
+ }
}
else if (region_type & ORT_ACC)
{
@@ -12075,6 +12138,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_DEFAULTMAP_NONE:
ctx->defaultmap[gdmk] = 0;
break;
+ case OMP_CLAUSE_DEFAULTMAP_PRESENT:
+ ctx->defaultmap[gdmk] = GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
+ break;
case OMP_CLAUSE_DEFAULTMAP_DEFAULT:
switch (gdmk)
{
@@ -12683,6 +12749,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
case GOVD_MAP_FORCE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
+ case GOVD_MAP_FORCE_PRESENT | GOVD_MAP_ALLOC_ONLY:
+ kind = GOMP_MAP_PRESENT_ALLOC;
+ break;
case GOVD_DEVICEPTR:
kind = GOMP_MAP_FORCE_DEVICEPTR;
break;
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index e746d458432..b485e21f23b 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1831,6 +1831,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TOFROM
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_PRESENT_TO
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_PRESENT_FROM
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_PRESENT_TOFROM
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable
@@ -13775,6 +13778,14 @@ 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_ALWAYS_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_STRUCT:
@@ -14503,6 +14514,13 @@ 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:
case GOMP_MAP_RELEASE:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
@@ -14545,11 +14563,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tkind_zero = tkind;
break;
case OMP_CLAUSE_TO:
- tkind = GOMP_MAP_TO;
+ tkind
+ = (OMP_CLAUSE_MOTION_PRESENT (c)
+ ? GOMP_MAP_PRESENT_TO : GOMP_MAP_TO);
tkind_zero = tkind;
break;
case OMP_CLAUSE_FROM:
- tkind = GOMP_MAP_FROM;
+ tkind
+ = (OMP_CLAUSE_MOTION_PRESENT (c)
+ ? GOMP_MAP_PRESENT_FROM : GOMP_MAP_FROM);
tkind_zero = tkind;
break;
default:
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 583294a7dcc..99f5b896493 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,25 @@
+2023-06-07 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from mainline:
+ 2023-06-06 Kwok Cheung Yeung <kcy@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+
+
+ * c-c++-common/gomp/map-6.c: Update dg-error, extend to test for
+ duplicated 'present' and extend scan-dump tests for 'present'.
+ * gfortran.dg/gomp/defaultmap-1.f90: Update dg-error.
+ * gfortran.dg/gomp/map-7.f90: Extend parse and dump test for
+ 'present'.
+ * gfortran.dg/gomp/map-8.f90: Extend for duplicate 'present'
+ modifier checking.
+ * c-c++-common/gomp/defaultmap-4.c: New test.
+ * c-c++-common/gomp/map-9.c: New test.
+ * c-c++-common/gomp/target-update-1.c: New test.
+ * gfortran.dg/gomp/defaultmap-8.f90: New test.
+ * gfortran.dg/gomp/map-11.f90: New test.
+ * gfortran.dg/gomp/map-12.f90: New test.
+ * gfortran.dg/gomp/target-update-1.f90: New test.
+
2023-06-07 Tobias Burnus <tobias@codesourcery.com>
Reverted:
diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
index f1ce7399310..bbc8fb4e32b 100644
--- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
@@ -15,7 +15,7 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
bar (p);
#pragma omp target map (p) , map (p[0])
bar (p);
- #pragma omp target map (q) map (q)
+ #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
bar (&q);
#pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */
bar (p);
diff --git a/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c
new file mode 100644
index 00000000000..1afff7ea38f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 1000
+
+void
+foo (void)
+{
+ int a[N], b[N], c[N];
+
+ /* Should generate implicit 'map(present, alloc)' clauses. */
+ #pragma omp target defaultmap (present: aggregate)
+ for (int i = 0; i < N; i++)
+ c[i] = a[i] + b[i];
+
+ /* Should generate implicit 'map(present, alloc)' clauses,
+ and they should go before other non-present clauses. */
+ #pragma omp target map(from: c) defaultmap (present: aggregate)
+ for (int i = 0; i < N; i++)
+ 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" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
index 6ee59714847..8c5b7f7cca1 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -13,10 +13,10 @@ 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' or 'close'" } */
+ #pragma omp target map (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 "'#pragma omp target' with modifier other than 'always' or 'close'" } */
+ #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) /* { dg-error "'close' undeclared" "" { target c } } */
@@ -80,9 +80,56 @@ foo (void)
#pragma omp target map (close close to:a) /* { dg-error "too many 'close' modifiers" } */
;
- #pragma omp target map (always to : a) map (close to : b)
+ #pragma omp target map (present , present , to:a) /* { dg-error "too many 'present' modifiers" } */
;
+ #pragma omp target map (present present , to:a) /* { dg-error "too many 'present' modifiers" } */
+ ;
+
+ #pragma omp target map (present , present to:a) /* { dg-error "too many 'present' modifiers" } */
+ ;
+
+ #pragma omp target map (present present to:a) /* { dg-error "too many 'present' modifiers" } */
+ ;
+
+ #pragma omp target map (always to : a) map (close to : b) map (present,tofrom : b1) map(always,present,alloc : b2) map(close, from: b3)
+ ;
+
+ #pragma omp target data map(tofrom:b1)
+ ;
+ #pragma omp target data map(close,tofrom:b1)
+ ;
+ #pragma omp target data map(close always,tofrom:b1)
+ ;
+ #pragma omp target data map(close always,tofrom:b1)
+ ;
+ #pragma omp target data map(close present,tofrom:b1)
+ ;
+ #pragma omp target data map(close present,tofrom:b1)
+ ;
+ #pragma omp target data map(always close present,tofrom:b1)
+ ;
+ #pragma omp target data map(always close present,tofrom:b1)
+ ;
+
+ #pragma omp target enter data map(alloc: a) map(to:b) map(tofrom:b1)
+ #pragma omp target enter data map(close, alloc: a) map(close,to:b) map(close,tofrom:b1)
+ #pragma omp target enter data map(always,alloc: a) map(always,to:b) map(close always,tofrom:b1)
+ #pragma omp target enter data map(always,close,alloc: a) map(close,always,to:b) map(close always,tofrom:b1)
+ #pragma omp target enter data map(present,alloc: a) map(present,to:b) map(close present,tofrom:b1)
+ #pragma omp target enter data map(present,close,alloc: a) map(close,present,to:b) map(close present,tofrom:b1)
+ #pragma omp target enter data map(present,always,alloc: a) map(always,present,to:b) map(always close present,tofrom:b1)
+ #pragma omp target enter data map(present,always,close,alloc: a) map(close,present,always,to:b) map(always close present,tofrom:b1)
+
+ #pragma omp target exit data map(delete: a) map(release:b) map(from:b1)
+ #pragma omp target exit data map(close,delete: a) map(close,release:b) map(close,from:b1)
+ #pragma omp target exit data map(always,delete: a) map(always,release:b) map(close always,from:b1)
+ #pragma omp target exit data map(always,close,delete: a) map(close,always,release:b) map(close always,from:b1)
+ #pragma omp target exit data map(present,delete: a) map(present,release:b) map(close present,from:b1)
+ #pragma omp target exit data map(present,close,delete: a) map(close,present,release:b) map(close present,from:b1)
+ #pragma omp target exit data map(present,always,delete: a) map(always,present,release:b) map(always close present,from:b1)
+ #pragma omp target exit data map(present,always,close,delete: a) map(close,present,always,release:b) map(always close present,from:b1)
+
int close = 0;
#pragma omp target map (close)
;
@@ -133,3 +180,20 @@ foo (void)
/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b5" "original" } } */
/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b6" "original" } } */
/* { dg-final { scan-tree-dump "pragma omp target map\\(always,to:b7\\) map\\(always,to:close\\) map\\(always,to:always\\)" "original" } } */
+
+/* Note: 'always,alloc' is the same as 'alloc'; hence, there is no 'always' for 'b'. Additionally, 'close' is ignored. */
+
+/* { dg-final { scan-tree-dump "#pragma omp target map\\(from:b3\\) map\\(present,alloc:b2\\) map\\(present,tofrom:b1\\) map\\(to:b\\) map\\(always,to:a\\)" "original" } } */
+
+/* { 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\\(to:b1\\) map\\(to:b\\) map\\(alloc:a\\)\[\r\n\]" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(always,to:b1\\) map\\(always,to:b\\) map\\(alloc:a\\)\[\r\n\]" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(present,to:b1\\) map\\(present,to:b\\) map\\(present,alloc:a\\)\[\r\n\]" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(always,present,to:b1\\) map\\(always,present,to:b\\) map\\(present,alloc:a\\)\[\r\n\]" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(from:b1\\) map\\(release:b\\) map\\(delete:a\\)\[\r\n\]" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,from:b1\\) map\\(release:b\\) map\\(delete:a\\)\[\r\n\]" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(present,from:b1\\) map\\(release:b\\) map\\(delete:a\\)\[\r\n\]" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(always,present,from:b1\\) map\\(release:b\\) map\\(delete:a\\)\[\r\n\]" 2 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/map-9.c b/gcc/testsuite/c-c++-common/gomp/map-9.c
new file mode 100644
index 00000000000..4b4bd6d2aa3
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-9.c
@@ -0,0 +1,32 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 1000
+
+void
+foo (void)
+{
+ int a[N], b[N], c[N];
+
+ /* Should be able to parse 'present' map modifier. */
+ #pragma omp target enter data map (present, to: a, b)
+
+ #pragma omp target data map (present, to: a, b) map (always, present, from: c)
+
+ #pragma omp target map (present, to: a, b) map (present, from: c)
+ for (int i = 0; i < N; i++)
+ c[i] = a[i] + b[i];
+
+ #pragma omp target exit data map (always, present, from: c)
+
+ /* Map clauses with 'present' modifier should go ahead of those without. */
+ #pragma omp target map (to: a) map (present, to: b) map (from: c)
+ for (int i = 0; i < N; i++)
+ 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 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" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-1.c b/gcc/testsuite/c-c++-common/gomp/target-update-1.c
new file mode 100644
index 00000000000..4408ab75255
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-1.c
@@ -0,0 +1,15 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 1000
+
+void
+foo (void)
+{
+ int a[N], b[N], c, d, e;
+
+ /* Should be able to parse present in to/from clauses of 'target update'. */
+ #pragma omp target update to(c) to(present: a) from(d) from(present: b) to(e)
+}
+
+/* { dg-final { scan-tree-dump "#pragma omp target update to\\(e \\\[len: \[0-9\]+\\\]\\) from\\(present:b \\\[len: \[0-9\]+\\\]\\) from\\(d \\\[len: \[0-9\]+\\\]\\) to\\(present:a \\\[len: \[0-9\]+\\\]\\) to\\(c \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90
index 299d971f23c..1f1b8528aef 100644
--- a/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90
@@ -2,7 +2,7 @@
implicit none
-!$omp target defaultmap(bar) ! { dg-error "25: Expected ALLOC, TO, FROM, TOFROM, FIRSTPRIVATE, NONE or DEFAULT" }
+!$omp target defaultmap(bar) ! { dg-error "25: Expected ALLOC, TO, FROM, TOFROM, FIRSTPRIVATE, PRESENT, NONE or DEFAULT" }
!$omp target defaultmap ( alloc: foo) ! { dg-error "34: Expected SCALAR, AGGREGATE, ALLOCATABLE or POINTER" }
diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90
new file mode 100644
index 00000000000..669a623f746
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90
@@ -0,0 +1,26 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ implicit none
+ integer, parameter :: N = 1000
+ integer :: a(N), b(N), c(N), i
+
+ ! Should generate implicit 'map(present, alloc)' clauses.
+ !$omp target defaultmap (present: aggregate)
+ do i = 1, N
+ c(i) = a(i) + b(i)
+ end do
+ !$omp end target
+
+ ! Should generate implicit 'map(present, alloc)' clauses,
+ ! and they should go before other non-present clauses.
+ !$omp target map(from: c) defaultmap (present: aggregate)
+ do i = 1, N
+ c(i) = a(i) + b(i)
+ end do
+ !$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" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-11.f90 b/gcc/testsuite/gfortran.dg/gomp/map-11.f90
new file mode 100644
index 00000000000..9eb956f1aa1
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-11.f90
@@ -0,0 +1,34 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ implicit none
+ integer, parameter :: N = 1000
+ integer :: a(N), b(N), c(N), i
+
+ ! Should be able to parse 'present' map modifier.
+ !$omp target enter data map (present, to: a, b)
+
+ !$omp target data map (present, to: a, b) map (always, present, from: c)
+ !$omp target map (present, to: a, b) map (present, from: c)
+ do i = 1, N
+ c(i) = a(i) + b(i)
+ end do
+ !$omp end target
+ !$omp end target data
+
+ !$omp target exit data map (always, present, from: c)
+
+ ! Map clauses with 'present' modifier should go ahead of those without.
+ !$omp target map (to: a) map (present, to: b) map (from: c)
+ do i = 1, N
+ c(i) = a(i) + b(i)
+ end do
+ !$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 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" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-12.f90 b/gcc/testsuite/gfortran.dg/gomp/map-12.f90
new file mode 100644
index 00000000000..74bd01f5f5a
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-12.f90
@@ -0,0 +1,67 @@
+! { dg-additional-options "-fdump-tree-omplower" }
+
+subroutine foo
+ implicit none
+ integer :: a, b, b1
+
+ !$omp target data map(tofrom:b1)
+ block; end block
+ !$omp target data map(close,tofrom:b1)
+ block; end block
+ !$omp target data map(close always,tofrom:b1)
+ block; end block
+ !$omp target data map(close always,tofrom:b1)
+ block; end block
+ !$omp target data map(close present,tofrom:b1)
+ block; end block
+ !$omp target data map(close present,tofrom:b1)
+ block; end block
+ !$omp target data map(always close present,tofrom:b1)
+ block; end block
+ !$omp target data map(always close present,tofrom:b1)
+ block; end block
+
+ !$omp target enter data map(alloc: a) map(to:b) map(tofrom:b1)
+ !$omp target enter data map(close, alloc: a) map(close,to:b) map(close,tofrom:b1)
+ !$omp target enter data map(always,alloc: a) map(always,to:b) map(close always,tofrom:b1)
+ !$omp target enter data map(always,close,alloc: a) map(close,always,to:b) map(close always,tofrom:b1)
+ !$omp target enter data map(present,alloc: a) map(present,to:b) map(close present,tofrom:b1)
+ !$omp target enter data map(present,close,alloc: a) map(close,present,to:b) map(close present,tofrom:b1)
+ !$omp target enter data map(present,always,alloc: a) map(always,present,to:b) map(always close present,tofrom:b1)
+ !$omp target enter data map(present,always,close,alloc: a) map(close,present,always,to:b) map(always close present,tofrom:b1)
+
+ !$omp target exit data map(delete: a) map(release:b) map(from:b1)
+ !$omp target exit data map(close,delete: a) map(close,release:b) map(close,from:b1)
+ !$omp target exit data map(always,delete: a) map(always,release:b) map(close always,from:b1)
+ !$omp target exit data map(always,close,delete: a) map(close,always,release:b) map(close always,from:b1)
+ !$omp target exit data map(present,delete: a) map(present,release:b) map(close present,from:b1)
+ !$omp target exit data map(present,close,delete: a) map(close,present,release:b) map(close present,from:b1)
+ !$omp target exit data map(present,always,delete: a) map(always,present,release:b) map(always close present,from:b1)
+ !$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 \\\[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\\(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 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\\(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/map-7.f90 b/gcc/testsuite/gfortran.dg/gomp/map-7.f90
index 009c6d49547..317090acb50 100644
--- a/gcc/testsuite/gfortran.dg/gomp/map-7.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-7.f90
@@ -2,7 +2,7 @@
implicit none
-integer :: a, b, close, always, to
+integer :: a, b, close, always, to, present
!$omp target map(close)
!$omp end target
@@ -10,17 +10,43 @@ integer :: a, b, close, always, to
!$omp target map(always)
!$omp end target
+!$omp target map(present)
+!$omp end target
+
!$omp target map(always, close)
!$omp end target
+!$omp target map(always, close, present)
+!$omp end target
+
!$omp target map(always, close, to : always, close, a)
!$omp end target
+!$omp target map(always, close, present, to : always, close, present, a)
+!$omp end target
+
+
!$omp target map(to, always, close)
!$omp end target
+!$omp target map(present, to, always, close)
+!$omp end target
+
+!$omp target map ( present , from : present) map(close , alloc : close) , map ( always, tofrom: always )
+!$omp end target
+
end
! { dg-final { scan-tree-dump-not "map\\(\[^\n\r)]*close\[^\n\r)]*to:" "original" } }
-! { dg-final { scan-tree-dump "#pragma omp target map\\(always,to:always\\) map\\(always,to:close\\) map\\(always,to:a\\)" "original" } }
! { dg-final { scan-tree-dump-not "map\\(\[^\n\r)]*close\[^\n\r)]*to:" "original" } }
+
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(tofrom:close\\)\[\n\r]" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(tofrom:always\\)\[\n\r]" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(tofrom:present\\)\[\n\r]" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(tofrom:always\\) map\\(tofrom:close\\)\[\n\r]" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(tofrom:always\\) map\\(tofrom:close\\) map\\(tofrom:present\\)\[\n\r]" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(always,to:always\\) map\\(always,to:close\\) map\\(always,to:a\\)\[\n\r]" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(always,present,to:always\\) map\\(always,present,to:close\\) map\\(always,present,to:present\\) map\\(always,present,to:a\\)\[\n\r]" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(tofrom:to\\) map\\(tofrom:always\\) map\\(tofrom:close\\)\[\n\r]" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(tofrom:present\\) map\\(tofrom:to\\) map\\(tofrom:always\\) map\\(tofrom:close\\)\[\n\r]" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target map\\(present,from:present\\) map\\(alloc:close\\) map\\(always,tofrom:always\\)\[\n\r]" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-8.f90 b/gcc/testsuite/gfortran.dg/gomp/map-8.f90
index 92b802c67ed..15ebdd68b95 100644
--- a/gcc/testsuite/gfortran.dg/gomp/map-8.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-8.f90
@@ -28,7 +28,18 @@ integer :: a
!$omp target map(close close to : a) ! { dg-error "too many 'close' modifiers" }
! !$omp end target
+!$omp target map(present present, to : a) ! { dg-error "too many 'present' modifiers" }
+! !$omp end target
+!$omp target map(present, present to : a) ! { dg-error "too many 'present' modifiers" }
+! !$omp end target
+!$omp target map(present present to : a) ! { dg-error "too many 'present' modifiers" }
+! !$omp end target
+
+
!$omp target map(close close always always to : a) ! { dg-error "too many 'always' modifiers" }
! !$omp end target
+!$omp target map(present close always present to : a) ! { dg-error "too many 'present' modifiers" }
+! !$omp end target
+
end
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90
new file mode 100644
index 00000000000..a9db2f1a39f
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90
@@ -0,0 +1,13 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ implicit none
+ integer, parameter :: N = 1000
+ integer :: a(N), b(N), c, d, e
+
+ ! Should be able to parse present in to/from clauses of 'target update'.
+ !$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\\(d \\\[len: \[0-9\]+\\\]\\) from\\(present:b \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 8d439c9c718..a1e7cbdd6ef 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -597,7 +597,8 @@ enum omp_clause_defaultmap_kind {
OMP_CLAUSE_DEFAULTMAP_NONE = 6 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1),
OMP_CLAUSE_DEFAULTMAP_DEFAULT
= 7 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1),
- OMP_CLAUSE_DEFAULTMAP_MASK = 7 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1)
+ OMP_CLAUSE_DEFAULTMAP_PRESENT = 8 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1),
+ OMP_CLAUSE_DEFAULTMAP_MASK = 15 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1)
};
enum omp_clause_bind_kind {
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index c11fb31ed68..77331c2a142 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1073,6 +1073,27 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, "attach_zero_length_array_section");
break;
+ case GOMP_MAP_PRESENT_ALLOC:
+ pp_string (pp, "present,alloc");
+ break;
+ case GOMP_MAP_PRESENT_TO:
+ pp_string (pp, "present,to");
+ break;
+ case GOMP_MAP_PRESENT_FROM:
+ pp_string (pp, "present,from");
+ break;
+ case GOMP_MAP_PRESENT_TOFROM:
+ pp_string (pp, "present,tofrom");
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ pp_string (pp, "always,present,to");
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
+ pp_string (pp, "always,present,from");
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ pp_string (pp, "always,present,tofrom");
+ break;
case GOMP_MAP_NONCONTIG_ARRAY_TO:
pp_string (pp, "to,noncontig_array");
break;
@@ -1154,12 +1175,16 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case OMP_CLAUSE_FROM:
pp_string (pp, "from(");
+ if (OMP_CLAUSE_MOTION_PRESENT (clause))
+ pp_string (pp, "present:");
dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
spc, flags, false);
goto print_clause_size;
case OMP_CLAUSE_TO:
pp_string (pp, "to(");
+ if (OMP_CLAUSE_MOTION_PRESENT (clause))
+ pp_string (pp, "present:");
dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
spc, flags, false);
goto print_clause_size;
@@ -1326,6 +1351,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case OMP_CLAUSE_DEFAULTMAP_NONE:
pp_string (pp, "none");
break;
+ case OMP_CLAUSE_DEFAULTMAP_PRESENT:
+ pp_string (pp, "present");
+ break;
case OMP_CLAUSE_DEFAULTMAP_DEFAULT:
pp_string (pp, "default");
break;
diff --git a/gcc/tree.h b/gcc/tree.h
index 9b2fffa061f..df1158ab72d 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1766,6 +1766,9 @@ class auto_suppress_location_wrappers
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind \
= (unsigned int) (MAP_KIND))
+#define OMP_CLAUSE_MOTION_PRESENT(NODE) \
+ (OMP_CLAUSE_RANGE_CHECK (NODE, OMP_CLAUSE_FROM, OMP_CLAUSE_TO)->base.deprecated_flag)
+
/* Nonzero if this map clause is for array (rather than pointer) based array
section with zero bias. Both the non-decl OMP_CLAUSE_MAP and corresponding
OMP_CLAUSE_MAP with GOMP_MAP_POINTER are marked with this flag. */
diff --git a/include/ChangeLog.omp b/include/ChangeLog.omp
index 6a3b56a37d7..a2f85843c01 100644
--- a/include/ChangeLog.omp
+++ b/include/ChangeLog.omp
@@ -1,3 +1,20 @@
+2023-06-07 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from mainline:
+ 2023-06-06 Kwok Cheung Yeung <kcy@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+
+ * gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_5): New.
+ (GOMP_MAP_FLAG_FORCE): Redefine.
+ (GOMP_MAP_FLAG_PRESENT, GOMP_MAP_FLAG_ALWAYS_PRESENT): New.
+ (enum gomp_map_kind): Add map kinds with 'present' modifiers.
+ (GOMP_MAP_COPY_TO_P, GOMP_MAP_COPY_FROM_P): Evaluate to true for
+ map variants with 'present'
+ (GOMP_MAP_ALWAYS_TO_P, GOMP_MAP_ALWAYS_FROM_P): Evaluate to true
+ for map variants with 'always, present' modifiers.
+ (GOMP_MAP_ALWAYS): Redefine.
+ (GOMP_MAP_FORCE_P, GOMP_MAP_PRESENT_P): New.
+
2023-06-07 Tobias Burnus <tobias@codesourcery.com>
Reverted:
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 99684c927c7..7d7bfb425f4 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -42,6 +42,7 @@
#define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4)
#define GOMP_MAP_FLAG_SPECIAL_3 (1 << 5)
#define GOMP_MAP_FLAG_SPECIAL_4 (1 << 6)
+#define GOMP_MAP_FLAG_SPECIAL_5 (1 << 7)
#define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \
| GOMP_MAP_FLAG_SPECIAL_0)
#define GOMP_MAP_DEEP_COPY (GOMP_MAP_FLAG_SPECIAL_4 \
@@ -55,9 +56,14 @@
| GOMP_MAP_FLAG_SPECIAL_1 \
| GOMP_MAP_FLAG_SPECIAL_2 \
| GOMP_MAP_FLAG_SPECIAL_3 \
- | GOMP_MAP_FLAG_SPECIAL_4)
+ | GOMP_MAP_FLAG_SPECIAL_4 \
+ | GOMP_MAP_FLAG_SPECIAL_5)
/* Flag to force a specific behavior (or else, trigger a run-time error). */
-#define GOMP_MAP_FLAG_FORCE (1 << 7)
+#define GOMP_MAP_FLAG_FORCE (GOMP_MAP_FLAG_SPECIAL_5)
+#define GOMP_MAP_FLAG_PRESENT (GOMP_MAP_FLAG_SPECIAL_5 \
+ | GOMP_MAP_FLAG_SPECIAL_0)
+#define GOMP_MAP_FLAG_ALWAYS_PRESENT (GOMP_MAP_FLAG_SPECIAL_2 \
+ | GOMP_MAP_FLAG_PRESENT)
enum gomp_map_kind
{
@@ -130,6 +136,23 @@ 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),
+ /* Must already be present, unconditionally copy from device. */
+ GOMP_MAP_ALWAYS_PRESENT_FROM = (GOMP_MAP_FLAG_ALWAYS_PRESENT
+ | GOMP_MAP_FROM),
+ /* Must already be present, unconditionally copy to and from device. */
+ GOMP_MAP_ALWAYS_PRESENT_TOFROM = (GOMP_MAP_FLAG_ALWAYS_PRESENT
+ | GOMP_MAP_TOFROM),
/* Map a sparse struct; the address is the base of the structure, alignment
it's required alignment, and size is the number of adjacent entries
that belong to the struct. The adjacent entries should be sorted by
@@ -211,11 +234,11 @@ enum gomp_map_kind
};
#define GOMP_MAP_COPY_TO_P(X) \
- (!((X) & GOMP_MAP_FLAG_SPECIAL) \
+ ((!((X) & GOMP_MAP_FLAG_SPECIAL) || GOMP_MAP_PRESENT_P (X)) \
&& ((X) & GOMP_MAP_FLAG_TO))
#define GOMP_MAP_COPY_FROM_P(X) \
- (!((X) & GOMP_MAP_FLAG_SPECIAL) \
+ ((!((X) & GOMP_MAP_FLAG_SPECIAL) || GOMP_MAP_PRESENT_P (X)) \
&& ((X) & GOMP_MAP_FLAG_FROM))
#define GOMP_MAP_ALWAYS_POINTER_P(X) \
@@ -226,17 +249,27 @@ enum gomp_map_kind
|| (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
#define GOMP_MAP_ALWAYS_TO_P(X) \
- (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
+ (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM) \
+ || ((X) == GOMP_MAP_ALWAYS_PRESENT_TO) \
+ || ((X) == GOMP_MAP_ALWAYS_PRESENT_TOFROM))
#define GOMP_MAP_ALWAYS_FROM_P(X) \
- (((X) == GOMP_MAP_ALWAYS_FROM) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
+ (((X) == GOMP_MAP_ALWAYS_FROM) || ((X) == GOMP_MAP_ALWAYS_TOFROM) \
+ || ((X) == GOMP_MAP_ALWAYS_PRESENT_FROM) \
+ || ((X) == GOMP_MAP_ALWAYS_PRESENT_TOFROM))
#define GOMP_MAP_ALWAYS_P(X) \
- (GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM))
+ (GOMP_MAP_ALWAYS_TO_P (X) || GOMP_MAP_ALWAYS_FROM_P (X))
#define GOMP_MAP_IMPLICIT_P(X) \
(((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_IMPLICIT)
+#define GOMP_MAP_FORCE_P(X) \
+ (((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)
+
#define GOMP_MAP_NONCONTIG_ARRAY_P(X) \
((X) & GOMP_MAP_NONCONTIG_ARRAY)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 6a2967f9385..0c0d1bd60e9 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,23 @@
+2023-06-07 Tobias Burnus <tobias@codesourcery.com>
+
+ Backported from mainline:
+ 2023-06-06 Kwok Cheung Yeung <kcy@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+
+ * libgomp.texi (OpenMP 5.1 Impl. status): Set 'present' support for
+ defaultmap to 'Y', add 'Y' entry for 'present' on to/from/map clauses.
+ * target.c (gomp_to_device_kind_p): Add map kinds with 'present'
+ modifier.
+ (gomp_map_vars_existing): Use new GOMP_MAP_FORCE_P macro.
+ (gomp_map_vars_internal, gomp_update, gomp_target_rev):
+ Emit runtime error if memory region not present.
+ * testsuite/libgomp.c-c++-common/target-present-1.c: New test.
+ * testsuite/libgomp.c-c++-common/target-present-2.c: New test.
+ * testsuite/libgomp.c-c++-common/target-present-3.c: New test.
+ * testsuite/libgomp.fortran/target-present-1.f90: New test.
+ * testsuite/libgomp.fortran/target-present-2.f90: New test.
+ * testsuite/libgomp.fortran/target-present-3.f90: New test.
+
2023-06-07 Tobias Burnus <tobias@codesourcery.com>
Reverted:
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index e9e74ce8ddb..c6f1ec42e5e 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -310,7 +310,7 @@ The OpenMP 4.5 specification is fully supported.
@item @code{inoutset} argument to the @code{depend} clause @tab Y @tab
@item @code{private} and @code{firstprivate} argument to @code{default}
clause in C and C++ @tab Y @tab
-@item @code{present} argument to @code{defaultmap} clause @tab N @tab
+@item @code{present} argument to @code{defaultmap} clause @tab Y @tab
@item @code{omp_set_num_teams}, @code{omp_set_teams_thread_limit},
@code{omp_get_max_teams}, @code{omp_get_teams_thread_limit} runtime
routines @tab Y @tab
@@ -352,6 +352,8 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab
@item Optional comma between directive and clause in the @code{#pragma} form @tab Y @tab
@item @code{indirect} clause in @code{declare target} @tab N @tab
@item @code{device_type(nohost)}/@code{device_type(host)} for variables @tab N @tab
+@item @code{present} modifier to the @code{map}, @code{to} and @code{from}
+ clauses @tab Y @tab
@end multitable
diff --git a/libgomp/target.c b/libgomp/target.c
index 18e174d3029..89f36f22997 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -426,6 +426,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:
return false;
default:
return true;
@@ -746,7 +748,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
else
tgt_var->length = newn->host_end - newn->host_start;
- if ((kind & GOMP_MAP_FLAG_FORCE)
+ if (GOMP_MAP_FORCE_P (kind)
/* For implicit maps, old contained in new is valid. */
|| !(implicit_subset
/* Otherwise, new contained inside old is considered valid. */
@@ -1991,6 +1993,20 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
#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,
@@ -2602,6 +2618,20 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
}
}
+ else
+ {
+ int kind = get_kind (short_mapkind, kinds, i);
+
+ if (GOMP_MAP_PRESENT_P (kind))
+ {
+ /* 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 *) hostaddrs[i], devicep->target_id);
+ }
+ }
}
gomp_mutex_unlock (&devicep->lock);
}
@@ -3998,7 +4028,8 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
case GOMP_MAP_DELETE:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION:
- /* Assume it is present; look it up - but ignore otherwise. */
+ /* Assume it is present; look it up - but ignore unless the
+ present clause is there. */
case GOMP_MAP_ALLOC:
case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_ALLOC:
@@ -4010,6 +4041,12 @@ 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_ALWAYS_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
cdata[i].devaddr = devaddrs[i];
bool zero_len = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
@@ -4030,7 +4067,23 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
devaddrs[i] + sizes[i], zero_len);
cdata[i].present = n2 != NULL;
}
- if (!cdata[i].present
+ if (!cdata[i].present && GOMP_MAP_PRESENT_P (kind))
+ {
+ gomp_mutex_unlock (&devicep->lock);
+#ifdef HAVE_INTTYPES_H
+ gomp_fatal ("present clause: no corresponding data on "
+ "parent device at %p with size %"PRIu64,
+ (void *) (uintptr_t) devaddrs[i],
+ (uint64_t) sizes[i]);
+#else
+ gomp_fatal ("present clause: no corresponding data on "
+ "parent device at %p with size %lu",
+ (void *) (uintptr_t) devaddrs[i],
+ (unsigned long) sizes[i]);
+#endif
+ break;
+ }
+ else if (!cdata[i].present
&& kind != GOMP_MAP_DELETE
&& kind != GOMP_MAP_RELEASE
&& kind != GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION)
@@ -4059,8 +4112,7 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
&& (kind == GOMP_MAP_TO || kind == GOMP_MAP_TOFROM))
|| kind == GOMP_MAP_FORCE_TO
|| kind == GOMP_MAP_FORCE_TOFROM
- || kind == GOMP_MAP_ALWAYS_TO
- || kind == GOMP_MAP_ALWAYS_TOFROM)
+ || GOMP_MAP_ALWAYS_TO_P (kind))
{
gomp_copy_dev2host (devicep, aq,
(void *) (uintptr_t) devaddrs[i],
@@ -4285,6 +4337,10 @@ 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;
/* FALLTHRU */
case GOMP_MAP_FROM:
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
new file mode 100644
index 00000000000..12f154c91a8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
@@ -0,0 +1,27 @@
+/* { dg-do run { target offload_device } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#define N 100
+
+int main (void)
+{
+ int a[N], b[N], c[N];
+
+ for (int i = 0; i < N; i++) {
+ a[i] = i * 2;
+ b[i] = i * 3 + 1;
+ }
+
+ #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)
+ for (int i = 0; i < N; i++)
+ c[i] = a[i];
+
+ /* 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\]+\\\)" } */
+ #pragma omp target map (present, to: b)
+ for (int i = 0; i < N; i++)
+ c[i] += b[i];
+ #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
new file mode 100644
index 00000000000..d4debbab10b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
@@ -0,0 +1,27 @@
+/* { dg-do run { target offload_device } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#define N 100
+
+int main (void)
+{
+ int a[N], b[N], c[N];
+
+ for (int i = 0; i < N; i++) {
+ a[i] = i * 2;
+ b[i] = i * 3 + 1;
+ }
+
+ #pragma omp target enter data map (alloc: a, c)
+ /* a has already been allocated, so this should be okay. */
+ #pragma omp target defaultmap (present)
+ for (int i = 0; i < N; i++)
+ c[i] = a[i];
+
+ /* 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\]+\\\)" } */
+ #pragma omp target defaultmap (present)
+ for (int i = 0; i < N; i++)
+ c[i] += b[i];
+ #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
new file mode 100644
index 00000000000..9d8d8f8a335
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
@@ -0,0 +1,27 @@
+/* { dg-do run { target offload_device } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#include <stdio.h>
+
+#define N 100
+
+int main (void)
+{
+ int a[N], b[N], c[N];
+
+ for (int i = 0; i < N; i++) {
+ a[i] = i * 2;
+ b[i] = i * 3 + 1;
+ }
+
+ #pragma omp target enter data map (alloc: a, c)
+
+ /* This should work as a has already been allocated. */
+ #pragma omp target update to (present: a)
+
+ /* 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\]+\\\)" } */
+ #pragma omp target update to (present: b)
+
+ #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-1.f90 b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
new file mode 100644
index 00000000000..349dcb118b2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
@@ -0,0 +1,30 @@
+! { dg-do run { target offload_device } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+ implicit none
+ integer, parameter :: N = 100
+ integer :: a(N), b(N), c(N), i
+
+ do i = 1, N
+ a(i) = i * 2
+ b(i) = i * 3 + 1
+ end do
+
+ !$omp target enter data map (alloc: a)
+ ! a has already been allocated, so this should be okay.
+ !$omp target map (present, to: a)
+ do i = 1, N
+ c(i) = a(i)
+ end do
+ !$omp end target
+
+ ! 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\]+\\\)" }
+ !$omp target map (present, to: b)
+ do i = 1, N
+ c(i) = c(i) + b(i)
+ end do
+ !$omp end target
+ !$omp target exit data map (from: c)
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-2.f90 b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
new file mode 100644
index 00000000000..07e79d1b576
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
@@ -0,0 +1,30 @@
+! { dg-do run { target offload_device } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+ implicit none
+ integer, parameter :: N = 100
+ integer :: a(N), b(N), c(N), i
+
+ do i = 1, N
+ a(i) = i * 2
+ b(i) = i * 3 + 1
+ end do
+
+ !$omp target enter data map (alloc: a)
+ ! a has already been allocated, so this should be okay.
+ !$omp target defaultmap (present)
+ do i = 1, N
+ c(i) = a(i)
+ end do
+ !$omp end target
+
+ ! 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\]+\\\)" }
+ !$omp target defaultmap (present)
+ do i = 1, N
+ c(i) = c(i) + b(i)
+ end do
+ !$omp end target
+!$omp target exit data map (from: c)
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-3.f90 b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
new file mode 100644
index 00000000000..a2709eb6f17
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
@@ -0,0 +1,22 @@
+! { dg-do run { target offload_device } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+ implicit none
+ integer, parameter :: N = 100
+ integer :: a(N), b(N), c(N), i
+
+ do i = 1, N
+ a(i) = i * 2
+ b(i) = i * 3 + 1
+ 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 fail as b has not been allocated.
+ ! { dg-output "libgomp: present clause: not present on the device \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" }
+ !$omp target update to (present: b)
+ !$omp target exit data map (from: c)
+end program
^ permalink raw reply [flat|nested] 2+ messages in thread
* [gcc/devel/omp/gcc-13] openmp: Add support for the 'present' modifier
@ 2023-05-19 16:52 Kwok Yeung
0 siblings, 0 replies; 2+ messages in thread
From: Kwok Yeung @ 2023-05-19 16:52 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:6e3816fa47c56d87d21d90f67435a1ab1664d8db
commit 6e3816fa47c56d87d21d90f67435a1ab1664d8db
Author: Kwok Cheung Yeung <kcy@codesourcery.com>
Date: Fri Feb 3 13:04:21 2023 +0000
openmp: Add support for the 'present' modifier
This implements support for the OpenMP 5.1 'present' modifier, which can be
used in map clauses in the 'target', 'target data', 'target data enter' and
'target data exit' constructs, and in the 'to' and 'from' clauses of the
'target update' construct. It is also supported in defaultmap.
The modifier triggers a fatal runtime error if the data specified by the
clause is not already present on the target device. It can also be combined
with 'always' in map clauses.
2023-02-01 Kwok Cheung Yeung <kcy@codesourcery.com>
gcc/c/
* c-parser.cc (c_parser_omp_variable_list): Set default motion
modifier.
(c_parser_omp_var_list_parens): Add new parameter with default. Parse
'present' motion modifier and apply.
(c_parser_omp_clause_defaultmap): Parse 'present' in defaultmap.
(c_parser_omp_clause_map): Parse 'present' modifier in map clauses.
(c_parser_omp_clause_to): Allow use of 'present' in variable list.
(c_parser_omp_clause_from): Likewise.
(c_parser_omp_target_data): Allow map clauses with 'present'
modifiers.
(c_parser_omp_target_enter_data): Likewise.
(c_parser_omp_target_exit_data): Likewise.
(c_parser_omp_target): Likewise.
gcc/cp/
* parser.cc (cp_parser_omp_var_list_no_open): Add new parameter with
default. Parse 'present' motion modifier and apply.
(cp_parser_omp_clause_defaultmap): Parse 'present' in defaultmap.
(cp_parser_omp_clause_map): Parse 'present' modifier in map clauses.
(cp_parser_omp_all_clauses): Allow use of 'present' in 'to' and 'from'
clauses.
(cp_parser_omp_target_data): Allow map clauses with 'present'
modifiers.
(cp_parser_omp_target_enter_data): Likewise.
(cp_parser_omp_target_exit_data): Likewise.
* semantics.cc (finish_omp_target): Accept map clauses with 'present'
modifiers.
gcc/fortran/
* dump-parse-tree.cc (show_omp_namelist): Display 'present' map
modifier.
(show_omp_clauses): Display 'present' motion modifier for 'to'
and 'from' clauses.
* gfortran.h (enum gfc_omp_map_op): Add entries with 'present'
modifiers.
(enum gfc_omp_motion_modifier): New.
(struct gfc_omp_namelist): Add motion_modifier field.
* openmp.cc (gfc_match_omp_variable_list): Add new parameter with
default. Parse 'present' motion modifier and apply.
(gfc_match_omp_clauses): Parse 'present' in defaultmap, 'from'
clauses, 'map' clauses and 'to' clauses.
(resolve_omp_clauses): Allow 'present' modifiers on 'target',
'target data', 'target enter' and 'target exit' directives.
* trans-openmp.cc (gfc_omp_deep_map_kind_p): Handle map kinds with
'present' modifier.
(gfc_trans_omp_clauses): Apply 'present' modifiers to tree node for
'map', 'to' and 'from' clauses. Apply 'present' for defaultmap.
gcc/
* gimplify.cc (omp_notice_variable): Apply GOVD_MAP_ALLOC_ONLY flag
and defaultmap flags if the defaultmap has GOVD_MAP_FORCE_PRESENT flag
set.
(omp_get_attachment): Handle map clauses with 'present' modifier.
(omp_group_base): Likewise.
(gimplify_scan_omp_clauses): Reorder present maps to come first.
Set GOVD flags for present defaultmaps.
(gimplify_adjust_omp_clauses_1): Set map kind for present defaultmaps.
* omp-low.cc (scan_sharing_clauses): Handle 'always, present' map
clauses.
(lower_omp_target): Handle map clauses with 'present' modifier.
Handle 'to' and 'from' clauses with 'present'.
* tree-core.h (enum omp_clause_defaultmap_kind): Add
OMP_CLAUSE_DEFAULTMAP_PRESENT defaultmap kind.
(enum omp_clause_motion_modifier): New.
(struct tree_omp_clause): Add motion_modifier field.
* tree-pretty-print.cc (dump_omp_clause): Handle 'map', 'to' and
'from' clauses with 'present' modifier. Handle present defaultmap.
* tree.h (OMP_CLAUSE_MOTION_MODIFIER): New.
(OMP_CLAUSE_SET_MOTION_MODIFIER): New.
gcc/testsuite/
* c-c++-common/gomp/defaultmap-4.c: New.
* c-c++-common/gomp/map-6.c: Update expected error messages.
* c-c++-common/gomp/map-9.c: New.
* c-c++-common/gomp/target-update-1.c: New.
* gfortran.dg/gomp/defaultmap-1.f90: Update expected error messages.
* gfortran.dg/gomp/defaultmap-8.f90: New.
* gfortran.dg/gomp/map-10.f90: New.
* gfortran.dg/gomp/target-update-1.f90: New.
include/
* gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_5): New.
(GOMP_MAP_FLAG_FORCE): Redefine.
(GOMP_MAP_FLAG_PRESENT): New.
(GOMP_MAP_FLAG_ALWAYS_PRESENT): New.
(enum gomp_map_kind): Add map kinds with 'present' modifiers.
(GOMP_MAP_COPY_TO_P): Evaluate to true for map variants with 'present'
modifiers.
(GOMP_MAP_COPY_FROM_P): Likewise.
(GOMP_MAP_ALWAYS_TO_P): Evaluate to true for map variants with
'always, present' modifiers.
(GOMP_MAP_ALWAYS_FROM_P): Likewise.
(GOMP_MAP_ALWAYS): Redefine.
(GOMP_MAP_FORCE_P): New.
(GOMP_MAP_PRESENT_P): New.
libgomp/
* target.c (gomp_to_device_kind_p): Add map kinds with 'present'
modifier.
(gomp_map_vars_existing): Use new GOMP_MAP_FORCE_P macro.
(gomp_map_vars_internal): Emit runtime error if memory region not
present.
(gomp_update): Likewise.
(gomp_target_rev): Likewise.
* testsuite/libgomp.c-c++-common/target-present-1.c: New.
* testsuite/libgomp.c-c++-common/target-present-2.c: New.
* testsuite/libgomp.c-c++-common/target-present-3.c: New.
* testsuite/libgomp.fortran/target-present-1.f90: New.
* testsuite/libgomp.fortran/target-present-2.f90: New.
* testsuite/libgomp.fortran/target-present-3.f90: New.
Add 'present' map types to gfc_omp_deep_map_kind_p
Diff:
---
gcc/ChangeLog.omp | 23 +++++
gcc/c/ChangeLog.omp | 16 ++++
gcc/c/c-parser.cc | 104 +++++++++++++++++++--
gcc/cp/ChangeLog.omp | 15 +++
gcc/cp/parser.cc | 103 +++++++++++++++++---
gcc/cp/semantics.cc | 7 ++
gcc/fortran/ChangeLog.omp | 21 +++++
gcc/fortran/dump-parse-tree.cc | 15 +++
gcc/fortran/gfortran.h | 14 +++
gcc/fortran/openmp.cc | 77 +++++++++++++--
gcc/fortran/trans-openmp.cc | 37 ++++++++
gcc/gimplify.cc | 69 ++++++++++++++
gcc/omp-low.cc | 26 +++++-
gcc/testsuite/ChangeLog.omp | 11 +++
gcc/testsuite/c-c++-common/gomp/defaultmap-4.c | 24 +++++
gcc/testsuite/c-c++-common/gomp/map-6.c | 4 +-
gcc/testsuite/c-c++-common/gomp/map-9.c | 32 +++++++
gcc/testsuite/c-c++-common/gomp/target-update-1.c | 15 +++
gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90 | 2 +-
gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 | 26 ++++++
gcc/testsuite/gfortran.dg/gomp/map-10.f90 | 34 +++++++
gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 | 13 +++
gcc/tree-core.h | 9 +-
gcc/tree-pretty-print.cc | 28 ++++++
gcc/tree.h | 6 ++
include/ChangeLog.omp | 17 ++++
include/gomp-constants.h | 47 ++++++++--
libgomp/ChangeLog.omp | 16 ++++
libgomp/target.c | 66 ++++++++++++-
.../libgomp.c-c++-common/target-present-1.c | 27 ++++++
.../libgomp.c-c++-common/target-present-2.c | 27 ++++++
.../libgomp.c-c++-common/target-present-3.c | 27 ++++++
.../testsuite/libgomp.fortran/target-present-1.f90 | 30 ++++++
.../testsuite/libgomp.fortran/target-present-2.f90 | 30 ++++++
.../testsuite/libgomp.fortran/target-present-3.f90 | 22 +++++
35 files changed, 994 insertions(+), 46 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 08213522a24..6ff1d64cf00 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,26 @@
+2023-02-01 Kwok Cheung Yeung <kcy@codesourcery.com>
+
+ * gimplify.cc (omp_notice_variable): Apply GOVD_MAP_ALLOC_ONLY flag
+ and defaultmap flags if the defaultmap has GOVD_MAP_FORCE_PRESENT flag
+ set.
+ (omp_get_attachment): Handle map clauses with 'present' modifier.
+ (omp_group_base): Likewise.
+ (gimplify_scan_omp_clauses): Reorder present maps to come first.
+ Set GOVD flags for present defaultmaps.
+ (gimplify_adjust_omp_clauses_1): Set map kind for present defaultmaps.
+ * omp-low.cc (scan_sharing_clauses): Handle 'always, present' map
+ clauses.
+ (lower_omp_target): Handle map clauses with 'present' modifier.
+ Handle 'to' and 'from' clauses with 'present'.
+ * tree-core.h (enum omp_clause_defaultmap_kind): Add
+ OMP_CLAUSE_DEFAULTMAP_PRESENT defaultmap kind.
+ (enum omp_clause_motion_modifier): New.
+ (struct tree_omp_clause): Add motion_modifier field.
+ * tree-pretty-print.cc (dump_omp_clause): Handle 'map', 'to' and
+ 'from' clauses with 'present' modifier. Handle present defaultmap.
+ * tree.h (OMP_CLAUSE_MOTION_MODIFIER): New.
+ (OMP_CLAUSE_SET_MOTION_MODIFIER): New.
+
2023-01-20 Thomas Schwinge <thomas@codesourcery.com>
* collect2.cc (write_c_file_glob): Allow for
diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index d13c7feed0b..423e95a5171 100644
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,19 @@
+2023-02-01 Kwok Cheung Yeung <kcy@codesourcery.com>
+
+ * c-parser.cc (c_parser_omp_variable_list): Set default motion
+ modifier.
+ (c_parser_omp_var_list_parens): Add new parameter with default. Parse
+ 'present' motion modifier and apply.
+ (c_parser_omp_clause_defaultmap): Parse 'present' in defaultmap.
+ (c_parser_omp_clause_map): Parse 'present' modifier in map clauses.
+ (c_parser_omp_clause_to): Allow use of 'present' in variable list.
+ (c_parser_omp_clause_from): Likewise.
+ (c_parser_omp_target_data): Allow map clauses with 'present'
+ modifiers.
+ (c_parser_omp_target_enter_data): Likewise.
+ (c_parser_omp_target_exit_data): Likewise.
+ (c_parser_omp_target): Likewise.
+
2022-03-17 Kwok Cheung Yeung <kcy@codesourcery.com>
* c-typeck.cc (handle_omp_array_sections_1): Add check to ensure
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index cb2141140ff..64be7732ef9 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -14036,6 +14036,8 @@ c_parser_omp_variable_list (c_parser *parser,
tree u = build_omp_clause (clause_loc, kind);
OMP_CLAUSE_DECL (u) = t;
OMP_CLAUSE_CHAIN (u) = list;
+ if (kind == OMP_CLAUSE_FROM || kind == OMP_CLAUSE_TO)
+ OMP_CLAUSE_SET_MOTION_MODIFIER (u, OMP_CLAUSE_MOTION_NONE);
list = u;
}
}
@@ -14065,7 +14067,8 @@ static tree
c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
tree list,
enum c_omp_region_type ort = C_ORT_OMP,
- bool allow_deref = false)
+ bool allow_deref = false,
+ bool allow_present = false)
{
/* The clauses location. */
location_t loc = c_parser_peek_token (parser)->location;
@@ -14073,8 +14076,27 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
matching_parens parens;
if (parens.require_open (parser))
{
+ bool present = false;
+
+ if (allow_present)
+ {
+ c_token *token = c_parser_peek_token (parser);
+
+ if (token->type == CPP_NAME
+ && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0
+ && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+ {
+ present = true;
+ c_parser_consume_token (parser);
+ c_parser_consume_token (parser);
+ }
+ }
list = c_parser_omp_variable_list (parser, loc, kind, list, ort,
allow_deref);
+
+ if (present)
+ for (tree clause = list; clause; clause = OMP_CLAUSE_CHAIN (clause))
+ OMP_CLAUSE_SET_MOTION_MODIFIER (clause, OMP_CLAUSE_MOTION_PRESENT);
parens.skip_until_found_close (parser);
}
return list;
@@ -14960,6 +14982,13 @@ c_parser_omp_clause_defaultmap (c_parser *parser, tree list)
goto invalid_behavior;
break;
+ case 'p':
+ if (strcmp ("present", p) == 0)
+ behavior = OMP_CLAUSE_DEFAULTMAP_PRESENT;
+ else
+ goto invalid_behavior;
+ break;
+
case 't':
if (strcmp ("tofrom", p) == 0)
behavior = OMP_CLAUSE_DEFAULTMAP_TOFROM;
@@ -17338,6 +17367,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
int always_modifier = 0;
int close_modifier = 0;
+ int present_modifier = 0;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
c_token *tok = c_parser_peek_token (parser);
@@ -17369,11 +17399,21 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
}
close_modifier++;
}
+ else if (strcmp ("present", p) == 0)
+ {
+ if (present_modifier)
+ {
+ c_parser_error (parser, "too many %<present%> modifiers");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ present_modifier++;
+ }
else
{
c_parser_error (parser, "%<#pragma omp target%> with "
- "modifier other than %<always%> or "
- "%<close%> on %<map%> clause");
+ "modifier other than %<always%>, %<close%> "
+ "or %<present%> on %<map%> clause");
parens.skip_until_found_close (parser);
return list;
}
@@ -17385,14 +17425,25 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
&& c_parser_peek_2nd_token (parser)->type == CPP_COLON)
{
const char *p = IDENTIFIER_POINTER (c_parser_peek_token (parser)->value);
+ int always_present_modifier = always_modifier && present_modifier;
+
if (strcmp ("alloc", p) == 0)
- kind = GOMP_MAP_ALLOC;
+ kind = present_modifier ? GOMP_MAP_PRESENT_ALLOC : GOMP_MAP_ALLOC;
else if (strcmp ("to", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+ kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TO
+ : present_modifier ? GOMP_MAP_PRESENT_TO
+ : always_modifier ? GOMP_MAP_ALWAYS_TO
+ : GOMP_MAP_TO;
else if (strcmp ("from", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+ kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_FROM
+ : present_modifier ? GOMP_MAP_PRESENT_FROM
+ : always_modifier ? GOMP_MAP_ALWAYS_FROM
+ : GOMP_MAP_FROM;
else if (strcmp ("tofrom", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+ kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TOFROM
+ : present_modifier ? GOMP_MAP_PRESENT_TOFROM
+ : always_modifier ? GOMP_MAP_ALWAYS_TOFROM
+ : GOMP_MAP_TOFROM;
else if (strcmp ("release", p) == 0)
kind = GOMP_MAP_RELEASE;
else if (strcmp ("delete", p) == 0)
@@ -17654,7 +17705,7 @@ static tree
c_parser_omp_clause_to (c_parser *parser, tree list)
{
return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_TO, list,
- C_ORT_OMP, true);
+ C_ORT_OMP, true, true);
}
/* OpenMP 4.0:
@@ -17664,7 +17715,7 @@ static tree
c_parser_omp_clause_from (c_parser *parser, tree list)
{
return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_FROM, list,
- C_ORT_OMP, true);
+ C_ORT_OMP, true, true);
}
/* OpenMP 4.0:
@@ -21990,11 +22041,18 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
@@ -22140,7 +22198,10 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
map_seen = 3;
break;
case GOMP_MAP_TOFROM:
@@ -22151,6 +22212,14 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_TO);
map_seen = 3;
break;
+ case GOMP_MAP_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_TO);
+ map_seen = 3;
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_TO);
+ map_seen = 3;
+ break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
@@ -22238,6 +22307,8 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
{
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
map_seen = 3;
@@ -22250,6 +22321,14 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_FROM);
map_seen = 3;
break;
+ case GOMP_MAP_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_FROM);
+ map_seen = 3;
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_FROM);
+ map_seen = 3;
+ break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
@@ -22496,11 +22575,18 @@ check_clauses:
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH_DETACH:
diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index 42d655b2c70..09c143c0872 100644
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,18 @@
+2023-02-01 Kwok Cheung Yeung <kcy@codesourcery.com>
+
+ * parser.cc (cp_parser_omp_var_list_no_open): Add new parameter with
+ default. Parse 'present' motion modifier and apply.
+ (cp_parser_omp_clause_defaultmap): Parse 'present' in defaultmap.
+ (cp_parser_omp_clause_map): Parse 'present' modifier in map clauses.
+ (cp_parser_omp_all_clauses): Allow use of 'present' in 'to' and 'from'
+ clauses.
+ (cp_parser_omp_target_data): Allow map clauses with 'present'
+ modifiers.
+ (cp_parser_omp_target_enter_data): Likewise.
+ (cp_parser_omp_target_exit_data): Likewise.
+ * semantics.cc (finish_omp_target): Accept map clauses with 'present'
+ modifiers.
+
2022-03-17 Kwok Cheung Yeung <kcy@codesourcery.com>
* semantics.cc (handle_omp_array_sections_1): Add check to ensure
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index d72dc1e06b5..5690505b016 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -37702,11 +37702,33 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
static tree
cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
enum c_omp_region_type ort = C_ORT_OMP,
- bool allow_deref = false)
+ bool allow_deref = false, bool allow_present = false)
{
if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
- return cp_parser_omp_var_list_no_open (parser, kind, list, NULL, ort,
- allow_deref);
+ {
+ bool present = false;
+
+ if (allow_present)
+ {
+ cp_token *token = cp_lexer_peek_token (parser->lexer);
+
+ if (token->type == CPP_NAME
+ && strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0
+ && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+ {
+ present = true;
+ cp_lexer_consume_token (parser->lexer);
+ cp_lexer_consume_token (parser->lexer);
+ }
+ }
+
+ list = cp_parser_omp_var_list_no_open (parser, kind, list, NULL, ort,
+ allow_deref);
+
+ if (present)
+ for (tree clause = list; clause; clause = OMP_CLAUSE_CHAIN (clause))
+ OMP_CLAUSE_SET_MOTION_MODIFIER (clause, OMP_CLAUSE_MOTION_PRESENT);
+ }
return list;
}
@@ -38758,6 +38780,13 @@ cp_parser_omp_clause_defaultmap (cp_parser *parser, tree list,
goto invalid_behavior;
break;
+ case 'p':
+ if (strcmp ("present", p) == 0)
+ behavior = OMP_CLAUSE_DEFAULTMAP_PRESENT;
+ else
+ goto invalid_behavior;
+ break;
+
case 't':
if (strcmp ("tofrom", p) == 0)
behavior = OMP_CLAUSE_DEFAULTMAP_TOFROM;
@@ -40738,6 +40767,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
bool always_modifier = false;
bool close_modifier = false;
+ bool present_modifier = false;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
cp_token *tok = cp_lexer_peek_token (parser->lexer);
@@ -40774,11 +40804,24 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
}
close_modifier = true;
}
+ else if (strcmp ("present", p) == 0)
+ {
+ if (present_modifier)
+ {
+ cp_parser_error (parser, "too many %<present%> modifiers");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ present_modifier = true;
+ }
else
{
cp_parser_error (parser, "%<#pragma omp target%> with "
- "modifier other than %<always%> or "
- "%<close%> on %<map%> clause");
+ "modifier other than %<always%>, %<close%> "
+ "or %<present%> on %<map%> clause");
cp_parser_skip_to_closing_parenthesis (parser,
/*recovering=*/true,
/*or_comma=*/false,
@@ -40794,15 +40837,25 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
{
tree id = cp_lexer_peek_token (parser->lexer)->u.value;
const char *p = IDENTIFIER_POINTER (id);
+ int always_present_modifier = always_modifier && present_modifier;
if (strcmp ("alloc", p) == 0)
- kind = GOMP_MAP_ALLOC;
+ kind = present_modifier ? GOMP_MAP_PRESENT_ALLOC : GOMP_MAP_ALLOC;
else if (strcmp ("to", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO;
+ kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TO
+ : present_modifier ? GOMP_MAP_PRESENT_TO
+ : always_modifier ? GOMP_MAP_ALWAYS_TO
+ : GOMP_MAP_TO;
else if (strcmp ("from", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_FROM : GOMP_MAP_FROM;
+ kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_FROM
+ : present_modifier ? GOMP_MAP_PRESENT_FROM
+ : always_modifier ? GOMP_MAP_ALWAYS_FROM
+ : GOMP_MAP_FROM;
else if (strcmp ("tofrom", p) == 0)
- kind = always_modifier ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
+ kind = always_present_modifier ? GOMP_MAP_ALWAYS_PRESENT_TOFROM
+ : present_modifier ? GOMP_MAP_PRESENT_TOFROM
+ : always_modifier ? GOMP_MAP_ALWAYS_TOFROM
+ : GOMP_MAP_TOFROM;
else if (strcmp ("release", p) == 0)
kind = GOMP_MAP_RELEASE;
else
@@ -41580,12 +41633,12 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
}
else
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_TO, clauses,
- C_ORT_OMP, true);
+ C_ORT_OMP, true, true);
c_name = "to";
break;
case PRAGMA_OMP_CLAUSE_FROM:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FROM, clauses,
- C_ORT_OMP, true);
+ C_ORT_OMP, true, true);
c_name = "from";
break;
case PRAGMA_OMP_CLAUSE_UNIFORM:
@@ -45423,11 +45476,18 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
@@ -45530,7 +45590,10 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
map_seen = 3;
break;
case GOMP_MAP_TOFROM:
@@ -45541,6 +45604,14 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_TO);
map_seen = 3;
break;
+ case GOMP_MAP_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_TO);
+ map_seen = 3;
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_TO);
+ map_seen = 3;
+ break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
@@ -45633,6 +45704,8 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
{
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
map_seen = 3;
@@ -45645,6 +45718,14 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_FROM);
map_seen = 3;
break;
+ case GOMP_MAP_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_PRESENT_FROM);
+ map_seen = 3;
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (*pc, GOMP_MAP_ALWAYS_PRESENT_FROM);
+ map_seen = 3;
+ break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index adf33270fee..38f97431359 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -10187,11 +10187,18 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
{
case GOMP_MAP_TO:
case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
case GOMP_MAP_TOFROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_ALWAYS_POINTER:
diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index f22722c39d2..07f8f7e0e5e 100644
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,24 @@
+2023-02-01 Kwok Cheung Yeung <kcy@codesourcery.com>
+
+ * dump-parse-tree.cc (show_omp_namelist): Display 'present' map
+ modifier.
+ (show_omp_clauses): Display 'present' motion modifier for 'to'
+ and 'from' clauses.
+ * gfortran.h (enum gfc_omp_map_op): Add entries with 'present'
+ modifiers.
+ (enum gfc_omp_motion_modifier): New.
+ (struct gfc_omp_namelist): Add motion_modifier field.
+ * openmp.cc (gfc_match_omp_variable_list): Add new parameter with
+ default. Parse 'present' motion modifier and apply.
+ (gfc_match_omp_clauses): Parse 'present' in defaultmap, 'from'
+ clauses, 'map' clauses and 'to' clauses.
+ (resolve_omp_clauses): Allow 'present' modifiers on 'target',
+ 'target data', 'target enter' and 'target exit' directives.
+ * trans-openmp.cc (gfc_omp_deep_map_kind_p): Handle map kinds with
+ 'present' modifier.
+ (gfc_trans_omp_clauses): Apply 'present' modifiers to tree node for
+ 'map', 'to' and 'from' clauses. Apply 'present' for defaultmap.
+
2022-11-02 Tobias Burnus <tobias@codesourcery.com>
* trans-openmp.cc (gfc_trans_omp_clauses): Ensure DT struct-comp with
diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc
index c32d3c701a2..060981dba87 100644
--- a/gcc/fortran/dump-parse-tree.cc
+++ b/gcc/fortran/dump-parse-tree.cc
@@ -1453,9 +1453,20 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n)
case OMP_MAP_TO: fputs ("to:", dumpfile); break;
case OMP_MAP_FROM: fputs ("from:", dumpfile); break;
case OMP_MAP_TOFROM: fputs ("tofrom:", dumpfile); break;
+ case OMP_MAP_PRESENT_ALLOC: fputs ("present,alloc:", dumpfile); break;
+ case OMP_MAP_PRESENT_TO: fputs ("present,to:", dumpfile); break;
+ case OMP_MAP_PRESENT_FROM: fputs ("present,from:", dumpfile); break;
+ case OMP_MAP_PRESENT_TOFROM:
+ fputs ("present,tofrom:", dumpfile); break;
case OMP_MAP_ALWAYS_TO: fputs ("always,to:", dumpfile); break;
case OMP_MAP_ALWAYS_FROM: fputs ("always,from:", dumpfile); break;
case OMP_MAP_ALWAYS_TOFROM: fputs ("always,tofrom:", dumpfile); break;
+ case OMP_MAP_ALWAYS_PRESENT_TO:
+ fputs ("always,present,to:", dumpfile); break;
+ case OMP_MAP_ALWAYS_PRESENT_FROM:
+ fputs ("always,present,from:", dumpfile); break;
+ case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+ fputs ("always,present,tofrom:", dumpfile); break;
case OMP_MAP_DELETE: fputs ("delete:", dumpfile); break;
case OMP_MAP_RELEASE: fputs ("release:", dumpfile); break;
default: break;
@@ -1793,6 +1804,10 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
fputs ("inscan, ", dumpfile);
if (list_type == OMP_LIST_REDUCTION_TASK)
fputs ("task, ", dumpfile);
+ if ((list_type == OMP_LIST_TO || list_type == OMP_LIST_FROM)
+ && omp_clauses->lists[list_type]->u.motion_modifier
+ == OMP_MOTION_PRESENT)
+ fputs ("present:", dumpfile);
show_omp_namelist (list_type, omp_clauses->lists[list_type]);
fputc (')', dumpfile);
}
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index e16e01929fe..2e134440e17 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1308,6 +1308,13 @@ enum gfc_omp_map_op
OMP_MAP_ALWAYS_TO,
OMP_MAP_ALWAYS_FROM,
OMP_MAP_ALWAYS_TOFROM,
+ OMP_MAP_PRESENT_ALLOC,
+ OMP_MAP_PRESENT_TO,
+ OMP_MAP_PRESENT_FROM,
+ OMP_MAP_PRESENT_TOFROM,
+ OMP_MAP_ALWAYS_PRESENT_TO,
+ OMP_MAP_ALWAYS_PRESENT_FROM,
+ OMP_MAP_ALWAYS_PRESENT_TOFROM,
OMP_MAP_DECLARE_ALLOCATE,
OMP_MAP_DECLARE_DEALLOCATE
};
@@ -1343,6 +1350,12 @@ enum gfc_omp_linear_op
OMP_LINEAR_UVAL
};
+enum gfc_omp_motion_modifier
+{
+ OMP_MOTION_NONE,
+ OMP_MOTION_PRESENT
+};
+
/* For use in OpenMP clauses in case we need extra information
(aligned clause alignment, linear clause step, etc.). */
@@ -1361,6 +1374,7 @@ typedef struct gfc_omp_namelist
ENUM_BITFIELD (gfc_omp_linear_op) op:4;
bool old_modifier;
} linear;
+ gfc_omp_motion_modifier motion_modifier;
struct gfc_common_head *common;
bool lastprivate_conditional;
} u;
diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc
index 6d4f1f99477..b8005e12815 100644
--- a/gcc/fortran/openmp.cc
+++ b/gcc/fortran/openmp.cc
@@ -417,7 +417,8 @@ gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
gfc_omp_namelist ***headp = NULL,
bool allow_sections = false,
bool allow_derived = false,
- bool *has_all_memory = NULL)
+ bool *has_all_memory = NULL,
+ bool allow_motion_modifier = false)
{
gfc_omp_namelist *head, *tail, *p;
locus old_loc, cur_loc;
@@ -425,6 +426,7 @@ gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
gfc_symbol *sym;
match m;
gfc_symtree *st;
+ bool present = false;
head = tail = NULL;
@@ -460,6 +462,12 @@ gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
tail->where = cur_loc;
goto next_item;
}
+ else if (allow_motion_modifier && m == MATCH_YES && strcmp (n, "present") == 0
+ && gfc_match_char (':') == MATCH_YES)
+ {
+ present = true;
+ m = gfc_match_name (n);
+ }
if (m == MATCH_YES)
{
gfc_symtree *st;
@@ -560,6 +568,13 @@ gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
*list = head;
if (headp)
*headp = list;
+
+ if (present)
+ {
+ gfc_omp_namelist *n;
+ for (n = head; n; n = n->next)
+ n->u.motion_modifier = OMP_MOTION_PRESENT;
+ }
return MATCH_YES;
syntax:
@@ -2633,6 +2648,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
behavior = OMP_DEFAULTMAP_FROM;
else if (gfc_match ("firstprivate ") == MATCH_YES)
behavior = OMP_DEFAULTMAP_FIRSTPRIVATE;
+ else if (gfc_match ("present ") == MATCH_YES)
+ behavior = OMP_DEFAULTMAP_PRESENT;
else if (gfc_match ("none ") == MATCH_YES)
behavior = OMP_DEFAULTMAP_NONE;
else if (gfc_match ("default ") == MATCH_YES)
@@ -2640,7 +2657,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
else
{
gfc_error ("Expected ALLOC, TO, FROM, TOFROM, FIRSTPRIVATE, "
- "NONE or DEFAULT at %C");
+ "PRESENT, NONE or DEFAULT at %C");
break;
}
if (')' == gfc_peek_ascii_char ())
@@ -3066,7 +3083,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_FROM)
&& (gfc_match_omp_variable_list ("from (",
&c->lists[OMP_LIST_FROM], false,
- NULL, &head, true, true)
+ NULL, &head, true, true, NULL,
+ true)
== MATCH_YES))
continue;
break;
@@ -3423,6 +3441,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
locus old_loc2 = gfc_current_locus;
int always_modifier = 0;
int close_modifier = 0;
+ int present_modifier = 0;
locus second_always_locus = old_loc2;
locus second_close_locus = old_loc2;
@@ -3439,20 +3458,38 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if (close_modifier++ == 1)
second_close_locus = current_locus;
}
+ else if (gfc_match ("present ") == MATCH_YES)
+ {
+ if (present_modifier++ == 1)
+ second_close_locus = current_locus;
+ }
else
break;
gfc_match (", ");
}
gfc_omp_map_op map_op = OMP_MAP_TOFROM;
+ int always_present_modifier
+ = always_modifier && present_modifier;
+
if (gfc_match ("alloc : ") == MATCH_YES)
- map_op = OMP_MAP_ALLOC;
+ map_op = present_modifier ? OMP_MAP_PRESENT_ALLOC
+ : OMP_MAP_ALLOC;
else if (gfc_match ("tofrom : ") == MATCH_YES)
- map_op = always_modifier ? OMP_MAP_ALWAYS_TOFROM : OMP_MAP_TOFROM;
+ map_op = always_present_modifier ? OMP_MAP_ALWAYS_PRESENT_TOFROM
+ : present_modifier ? OMP_MAP_PRESENT_TOFROM
+ : always_modifier ? OMP_MAP_ALWAYS_TOFROM
+ : OMP_MAP_TOFROM;
else if (gfc_match ("to : ") == MATCH_YES)
- map_op = always_modifier ? OMP_MAP_ALWAYS_TO : OMP_MAP_TO;
+ map_op = always_present_modifier ? OMP_MAP_ALWAYS_PRESENT_TO
+ : present_modifier ? OMP_MAP_PRESENT_TO
+ : always_modifier ? OMP_MAP_ALWAYS_TO
+ : OMP_MAP_TO;
else if (gfc_match ("from : ") == MATCH_YES)
- map_op = always_modifier ? OMP_MAP_ALWAYS_FROM : OMP_MAP_FROM;
+ map_op = always_present_modifier ? OMP_MAP_ALWAYS_PRESENT_FROM
+ : present_modifier ? OMP_MAP_PRESENT_FROM
+ : always_modifier ? OMP_MAP_ALWAYS_FROM
+ : OMP_MAP_FROM;
else if (gfc_match ("release : ") == MATCH_YES)
map_op = OMP_MAP_RELEASE;
else if (gfc_match ("delete : ") == MATCH_YES)
@@ -4004,7 +4041,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
else if ((mask & OMP_CLAUSE_TO)
&& (gfc_match_omp_variable_list ("to (",
&c->lists[OMP_LIST_TO], false,
- NULL, &head, true, true)
+ NULL, &head, true, true, NULL,
+ true)
== MATCH_YES))
continue;
break;
@@ -8592,11 +8630,18 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
{
case OMP_MAP_TO:
case OMP_MAP_ALWAYS_TO:
+ case OMP_MAP_PRESENT_TO:
+ case OMP_MAP_ALWAYS_PRESENT_TO:
case OMP_MAP_FROM:
case OMP_MAP_ALWAYS_FROM:
+ case OMP_MAP_PRESENT_FROM:
+ case OMP_MAP_ALWAYS_PRESENT_FROM:
case OMP_MAP_TOFROM:
case OMP_MAP_ALWAYS_TOFROM:
+ case OMP_MAP_PRESENT_TOFROM:
+ case OMP_MAP_ALWAYS_PRESENT_TOFROM:
case OMP_MAP_ALLOC:
+ case OMP_MAP_PRESENT_ALLOC:
break;
default:
gfc_error ("TARGET%s with map-type other than TO, "
@@ -8612,6 +8657,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
{
case OMP_MAP_TO:
case OMP_MAP_ALWAYS_TO:
+ case OMP_MAP_PRESENT_TO:
+ case OMP_MAP_ALWAYS_PRESENT_TO:
case OMP_MAP_ALLOC:
break;
case OMP_MAP_TOFROM:
@@ -8620,6 +8667,12 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
case OMP_MAP_ALWAYS_TOFROM:
n->u.map_op = OMP_MAP_ALWAYS_TO;
break;
+ case OMP_MAP_PRESENT_TOFROM:
+ n->u.map_op = OMP_MAP_PRESENT_TO;
+ break;
+ case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+ n->u.map_op = OMP_MAP_ALWAYS_PRESENT_TO;
+ break;
default:
gfc_error ("TARGET ENTER DATA with map-type other "
"than TO, TOFROM or ALLOC on MAP clause "
@@ -8632,6 +8685,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
{
case OMP_MAP_FROM:
case OMP_MAP_ALWAYS_FROM:
+ case OMP_MAP_PRESENT_FROM:
+ case OMP_MAP_ALWAYS_PRESENT_FROM:
case OMP_MAP_RELEASE:
case OMP_MAP_DELETE:
break;
@@ -8641,6 +8696,12 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
case OMP_MAP_ALWAYS_TOFROM:
n->u.map_op = OMP_MAP_ALWAYS_FROM;
break;
+ case OMP_MAP_PRESENT_TOFROM:
+ n->u.map_op = OMP_MAP_PRESENT_FROM;
+ break;
+ case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+ n->u.map_op = OMP_MAP_ALWAYS_PRESENT_FROM;
+ break;
default:
gfc_error ("TARGET EXIT DATA with map-type other "
"than FROM, TOFROM, RELEASE, or DELETE on "
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index 1bc87757722..260c07c7e6d 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -2882,9 +2882,16 @@ gfc_omp_deep_map_kind_p (tree clause)
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
+ 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_FIRSTPRIVATE:
return true;
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_PRESENT_ALLOC:
case GOMP_MAP_POINTER:
case GOMP_MAP_TO_PSET:
case GOMP_MAP_FORCE_PRESENT:
@@ -4461,6 +4468,30 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
always_modifier = true;
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TOFROM);
break;
+ case OMP_MAP_PRESENT_ALLOC:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_ALLOC);
+ break;
+ case OMP_MAP_PRESENT_TO:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_TO);
+ break;
+ case OMP_MAP_PRESENT_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_FROM);
+ break;
+ case OMP_MAP_PRESENT_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_PRESENT_TOFROM);
+ break;
+ case OMP_MAP_ALWAYS_PRESENT_TO:
+ always_modifier = true;
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_PRESENT_TO);
+ break;
+ case OMP_MAP_ALWAYS_PRESENT_FROM:
+ always_modifier = true;
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_PRESENT_FROM);
+ break;
+ case OMP_MAP_ALWAYS_PRESENT_TOFROM:
+ always_modifier = true;
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_PRESENT_TOFROM);
+ break;
case OMP_MAP_RELEASE:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
break;
@@ -5208,6 +5239,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
}
+ if (n->u.motion_modifier == OMP_MOTION_PRESENT)
+ OMP_CLAUSE_SET_MOTION_MODIFIER (node,
+ OMP_CLAUSE_MOTION_PRESENT);
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
}
break;
@@ -5772,6 +5806,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_DEFAULTMAP_FIRSTPRIVATE:
behavior = OMP_CLAUSE_DEFAULTMAP_FIRSTPRIVATE;
break;
+ case OMP_DEFAULTMAP_PRESENT:
+ behavior = OMP_CLAUSE_DEFAULTMAP_PRESENT;
+ break;
case OMP_DEFAULTMAP_NONE: behavior = OMP_CLAUSE_DEFAULTMAP_NONE; break;
case OMP_DEFAULTMAP_DEFAULT:
behavior = OMP_CLAUSE_DEFAULTMAP_DEFAULT;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index ecb0a7a181e..d3e0711b8bd 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -7934,6 +7934,11 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
else if (ctx->defaultmap[gdmk]
& (GOVD_MAP_0LEN_ARRAY | GOVD_FIRSTPRIVATE))
nflags |= ctx->defaultmap[gdmk];
+ else if (ctx->defaultmap[gdmk] & GOVD_MAP_FORCE_PRESENT)
+ {
+ gcc_assert (ctx->defaultmap[gdmk] & GOVD_MAP);
+ nflags |= ctx->defaultmap[gdmk] | GOVD_MAP_ALLOC_ONLY;
+ }
else
{
gcc_assert (ctx->defaultmap[gdmk] & GOVD_MAP);
@@ -9109,6 +9114,13 @@ omp_get_attachment (omp_mapping_group *grp)
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_PRESENT_ALLOC:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
@@ -9340,6 +9352,13 @@ omp_group_base (omp_mapping_group *grp, unsigned int *chained,
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_PRESENT_ALLOC:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_PRESENT_TOFROM:
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ALLOC:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
@@ -10822,6 +10841,50 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
delete grpmap;
delete groups;
}
+
+ /* OpenMP map clauses with 'present' need to go in front of those
+ without. */
+ tree present_map_head = NULL;
+ tree *present_map_tail_p = &present_map_head;
+ tree *first_map_clause_p = NULL;
+
+ for (tree *c_p = list_p; *c_p; )
+ {
+ tree c = *c_p;
+ tree *next_c_p = &OMP_CLAUSE_CHAIN (c);
+
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+ {
+ if (!first_map_clause_p)
+ first_map_clause_p = c_p;
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_PRESENT_ALLOC:
+ case GOMP_MAP_PRESENT_FROM:
+ case GOMP_MAP_PRESENT_TO:
+ case GOMP_MAP_PRESENT_TOFROM:
+ next_c_p = c_p;
+ *c_p = OMP_CLAUSE_CHAIN (c);
+
+ OMP_CLAUSE_CHAIN (c) = NULL;
+ *present_map_tail_p = c;
+ present_map_tail_p = &OMP_CLAUSE_CHAIN (c);
+
+ break;
+
+ default:
+ break;
+ }
+ }
+
+ c_p = next_c_p;
+ }
+ if (first_map_clause_p && present_map_head)
+ {
+ tree next = *first_map_clause_p;
+ *first_map_clause_p = present_map_head;
+ *present_map_tail_p = next;
+ }
}
else if (region_type & ORT_ACC)
{
@@ -12084,6 +12147,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_DEFAULTMAP_NONE:
ctx->defaultmap[gdmk] = 0;
break;
+ case OMP_CLAUSE_DEFAULTMAP_PRESENT:
+ ctx->defaultmap[gdmk] = GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
+ break;
case OMP_CLAUSE_DEFAULTMAP_DEFAULT:
switch (gdmk)
{
@@ -12690,6 +12756,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
case GOVD_DEVICEPTR:
kind = GOMP_MAP_FORCE_DEVICEPTR;
break;
+ case GOVD_MAP_FORCE_PRESENT | GOVD_MAP_ALLOC_ONLY:
+ kind = GOMP_MAP_PRESENT_ALLOC;
+ break;
default:
gcc_unreachable ();
}
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index a63139b1b98..0f7b3cdc085 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1827,6 +1827,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TOFROM
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_PRESENT_TO
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_PRESENT_FROM
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_PRESENT_TOFROM
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable
@@ -13675,6 +13678,14 @@ 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_ALWAYS_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+
case GOMP_MAP_FIRSTPRIVATE_POINTER:
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_STRUCT:
@@ -14403,6 +14414,13 @@ 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:
case GOMP_MAP_RELEASE:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
@@ -14445,11 +14463,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tkind_zero = tkind;
break;
case OMP_CLAUSE_TO:
- tkind = GOMP_MAP_TO;
+ tkind
+ = OMP_CLAUSE_MOTION_MODIFIER (c) == OMP_CLAUSE_MOTION_PRESENT
+ ? GOMP_MAP_PRESENT_TO : GOMP_MAP_TO;
tkind_zero = tkind;
break;
case OMP_CLAUSE_FROM:
- tkind = GOMP_MAP_FROM;
+ tkind
+ = OMP_CLAUSE_MOTION_MODIFIER (c) == OMP_CLAUSE_MOTION_PRESENT
+ ? GOMP_MAP_PRESENT_FROM : GOMP_MAP_FROM;
tkind_zero = tkind;
break;
default:
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 4a09719770e..33385d9ad11 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,14 @@
+2023-02-01 Kwok Cheung Yeung <kcy@codesourcery.com>
+
+ * c-c++-common/gomp/defaultmap-4.c: New.
+ * c-c++-common/gomp/map-6.c: Update expected error messages.
+ * c-c++-common/gomp/map-9.c: New.
+ * c-c++-common/gomp/target-update-1.c: New.
+ * gfortran.dg/gomp/defaultmap-1.f90: Update expected error messages.
+ * gfortran.dg/gomp/defaultmap-8.f90: New.
+ * gfortran.dg/gomp/map-10.f90: New.
+ * gfortran.dg/gomp/target-update-1.f90: New.
+
2023-02-09 Thomas Schwinge <thomas@codesourcery.com>
* gfortran.dg/gomp/allocate-4.f90: Cut.
diff --git a/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c
new file mode 100644
index 00000000000..1afff7ea38f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 1000
+
+void
+foo (void)
+{
+ int a[N], b[N], c[N];
+
+ /* Should generate implicit 'map(present, alloc)' clauses. */
+ #pragma omp target defaultmap (present: aggregate)
+ for (int i = 0; i < N; i++)
+ c[i] = a[i] + b[i];
+
+ /* Should generate implicit 'map(present, alloc)' clauses,
+ and they should go before other non-present clauses. */
+ #pragma omp target map(from: c) defaultmap (present: aggregate)
+ for (int i = 0; i < N; i++)
+ 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" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
index 6ee59714847..b4683ddbabf 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -13,10 +13,10 @@ 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' or 'close'" } */
+ #pragma omp target map (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 "'#pragma omp target' with modifier other than 'always' or 'close'" } */
+ #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) /* { dg-error "'close' undeclared" "" { target c } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/map-9.c b/gcc/testsuite/c-c++-common/gomp/map-9.c
new file mode 100644
index 00000000000..4b4bd6d2aa3
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-9.c
@@ -0,0 +1,32 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 1000
+
+void
+foo (void)
+{
+ int a[N], b[N], c[N];
+
+ /* Should be able to parse 'present' map modifier. */
+ #pragma omp target enter data map (present, to: a, b)
+
+ #pragma omp target data map (present, to: a, b) map (always, present, from: c)
+
+ #pragma omp target map (present, to: a, b) map (present, from: c)
+ for (int i = 0; i < N; i++)
+ c[i] = a[i] + b[i];
+
+ #pragma omp target exit data map (always, present, from: c)
+
+ /* Map clauses with 'present' modifier should go ahead of those without. */
+ #pragma omp target map (to: a) map (present, to: b) map (from: c)
+ for (int i = 0; i < N; i++)
+ 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 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" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-1.c b/gcc/testsuite/c-c++-common/gomp/target-update-1.c
new file mode 100644
index 00000000000..0233fe5a7af
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-1.c
@@ -0,0 +1,15 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 1000
+
+void
+foo (void)
+{
+ int a[N], b[N];
+
+ /* Should be able to parse present in to/from clauses of 'target update'. */
+ #pragma omp target update to(present: a) from(present: b)
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target update from\\(present:b \\\[len: \[0-9\]+\\\]\\) to\\(present:a \\\[len: \[0-9\]+\\\]\\)" "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90
index 299d971f23c..1f1b8528aef 100644
--- a/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-1.f90
@@ -2,7 +2,7 @@
implicit none
-!$omp target defaultmap(bar) ! { dg-error "25: Expected ALLOC, TO, FROM, TOFROM, FIRSTPRIVATE, NONE or DEFAULT" }
+!$omp target defaultmap(bar) ! { dg-error "25: Expected ALLOC, TO, FROM, TOFROM, FIRSTPRIVATE, PRESENT, NONE or DEFAULT" }
!$omp target defaultmap ( alloc: foo) ! { dg-error "34: Expected SCALAR, AGGREGATE, ALLOCATABLE or POINTER" }
diff --git a/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90 b/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90
new file mode 100644
index 00000000000..669a623f746
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/defaultmap-8.f90
@@ -0,0 +1,26 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ implicit none
+ integer, parameter :: N = 1000
+ integer :: a(N), b(N), c(N), i
+
+ ! Should generate implicit 'map(present, alloc)' clauses.
+ !$omp target defaultmap (present: aggregate)
+ do i = 1, N
+ c(i) = a(i) + b(i)
+ end do
+ !$omp end target
+
+ ! Should generate implicit 'map(present, alloc)' clauses,
+ ! and they should go before other non-present clauses.
+ !$omp target map(from: c) defaultmap (present: aggregate)
+ do i = 1, N
+ c(i) = a(i) + b(i)
+ end do
+ !$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" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-10.f90 b/gcc/testsuite/gfortran.dg/gomp/map-10.f90
new file mode 100644
index 00000000000..cc87212f8d0
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/map-10.f90
@@ -0,0 +1,34 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ implicit none
+ integer, parameter :: N = 1000
+ integer :: a(N), b(N), c(N), i
+
+ ! Should be able to parse 'present' map modifier.
+ !$omp target enter data map (present, to: a, b)
+
+ !$omp target data map (present, to: a, b) map (always, present, from: c)
+ !$omp target map (present, to: a, b) map (present, from: c)
+ do i = 1, N
+ c(i) = a(i) + b(i)
+ end do
+ !$omp end target
+ !$omp end target data
+
+ !$omp target exit data map (always, present, from: c)
+
+ ! Map clauses with 'present' modifier should go ahead of those without.
+ !$omp target map (to: a) map (present, to: b) map (from: c)
+ do i = 1, N
+ c(i) = a(i) + b(i)
+ end do
+ !$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 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" } }
diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90
new file mode 100644
index 00000000000..a382b87f229
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target-update-1.f90
@@ -0,0 +1,13 @@
+! { dg-do compile }
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program main
+ implicit none
+ integer, parameter :: N = 1000
+ integer :: a(N), b(N), i
+
+ ! Should be able to parse present in to/from clauses of 'target update'.
+ !$omp target update to(present: a) from(present: b)
+end program
+
+! { dg-final { scan-tree-dump "pragma omp target update to\\(present:a \\\[len: \[0-9\]+\\\]\\) from\\(present:b \\\[len: \[0-9\]+\\\]\\)" "gimple" } }
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 37b408ee480..7b6d8803969 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -581,7 +581,8 @@ enum omp_clause_defaultmap_kind {
OMP_CLAUSE_DEFAULTMAP_NONE = 6 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1),
OMP_CLAUSE_DEFAULTMAP_DEFAULT
= 7 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1),
- OMP_CLAUSE_DEFAULTMAP_MASK = 7 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1)
+ OMP_CLAUSE_DEFAULTMAP_PRESENT = 8 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1),
+ OMP_CLAUSE_DEFAULTMAP_MASK = 15 * (OMP_CLAUSE_DEFAULTMAP_CATEGORY_MASK + 1)
};
enum omp_clause_bind_kind {
@@ -590,6 +591,11 @@ enum omp_clause_bind_kind {
OMP_CLAUSE_BIND_THREAD
};
+enum omp_clause_motion_modifier {
+ OMP_CLAUSE_MOTION_NONE,
+ OMP_CLAUSE_MOTION_PRESENT
+};
+
/* memory-order-clause on OpenMP atomic/flush constructs or
argument of atomic_default_mem_order clause. */
enum omp_memory_order {
@@ -1660,6 +1666,7 @@ struct GTY(()) tree_omp_clause {
enum omp_clause_defaultmap_kind defaultmap_kind;
enum omp_clause_bind_kind bind_kind;
enum omp_clause_device_type_kind device_type_kind;
+ enum omp_clause_motion_modifier motion_modifier;
} GTY ((skip)) subcode;
/* The gimplification of OMP_CLAUSE_REDUCTION_{INIT,MERGE} for omp-low's
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 555dc5616e9..9046934face 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1052,6 +1052,27 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_NONCONTIG_ARRAY_FORCE_PRESENT:
pp_string (pp, "force_present,noncontig_array");
break;
+ case GOMP_MAP_PRESENT_ALLOC:
+ pp_string (pp, "present,alloc");
+ break;
+ case GOMP_MAP_PRESENT_TO:
+ pp_string (pp, "present,to");
+ break;
+ case GOMP_MAP_PRESENT_FROM:
+ pp_string (pp, "present,from");
+ break;
+ case GOMP_MAP_PRESENT_TOFROM:
+ pp_string (pp, "present,tofrom");
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ pp_string (pp, "always,present,to");
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_FROM:
+ pp_string (pp, "always,present,from");
+ break;
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
+ pp_string (pp, "always,present,tofrom");
+ break;
default:
gcc_unreachable ();
}
@@ -1106,12 +1127,16 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case OMP_CLAUSE_FROM:
pp_string (pp, "from(");
+ if (OMP_CLAUSE_MOTION_MODIFIER (clause) == OMP_CLAUSE_MOTION_PRESENT)
+ pp_string (pp, "present:");
dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
spc, flags, false);
goto print_clause_size;
case OMP_CLAUSE_TO:
pp_string (pp, "to(");
+ if (OMP_CLAUSE_MOTION_MODIFIER (clause) == OMP_CLAUSE_MOTION_PRESENT)
+ pp_string (pp, "present:");
dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
spc, flags, false);
goto print_clause_size;
@@ -1278,6 +1303,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case OMP_CLAUSE_DEFAULTMAP_NONE:
pp_string (pp, "none");
break;
+ case OMP_CLAUSE_DEFAULTMAP_PRESENT:
+ pp_string (pp, "present");
+ break;
case OMP_CLAUSE_DEFAULTMAP_DEFAULT:
pp_string (pp, "default");
break;
diff --git a/gcc/tree.h b/gcc/tree.h
index e20f1e7978e..273376b3d3b 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1766,6 +1766,12 @@ class auto_suppress_location_wrappers
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->omp_clause.subcode.map_kind \
= (unsigned int) (MAP_KIND))
+#define OMP_CLAUSE_MOTION_MODIFIER(NODE) \
+ ((enum omp_clause_motion_modifier) OMP_CLAUSE_RANGE_CHECK (NODE, OMP_CLAUSE_FROM, OMP_CLAUSE_TO)->omp_clause.subcode.motion_modifier)
+#define OMP_CLAUSE_SET_MOTION_MODIFIER(NODE, MOTION_MODIFIER) \
+ (OMP_CLAUSE_RANGE_CHECK (NODE, OMP_CLAUSE_FROM, OMP_CLAUSE_TO)->omp_clause.subcode.motion_modifier \
+ = (MOTION_MODIFIER))
+
/* Nonzero if this map clause is for array (rather than pointer) based array
section with zero bias. Both the non-decl OMP_CLAUSE_MAP and corresponding
OMP_CLAUSE_MAP with GOMP_MAP_POINTER are marked with this flag. */
diff --git a/include/ChangeLog.omp b/include/ChangeLog.omp
index 261be6b587a..f904e7a2763 100644
--- a/include/ChangeLog.omp
+++ b/include/ChangeLog.omp
@@ -1,3 +1,20 @@
+2023-02-01 Kwok Cheung Yeung <kcy@codesourcery.com>
+
+ * gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_5): New.
+ (GOMP_MAP_FLAG_FORCE): Redefine.
+ (GOMP_MAP_FLAG_PRESENT): New.
+ (GOMP_MAP_FLAG_ALWAYS_PRESENT): New.
+ (enum gomp_map_kind): Add map kinds with 'present' modifiers.
+ (GOMP_MAP_COPY_TO_P): Evaluate to true for map variants with 'present'
+ modifiers.
+ (GOMP_MAP_COPY_FROM_P): Likewise.
+ (GOMP_MAP_ALWAYS_TO_P): Evaluate to true for map variants with
+ 'always, present' modifiers.
+ (GOMP_MAP_ALWAYS_FROM_P): Likewise.
+ (GOMP_MAP_ALWAYS): Redefine.
+ (GOMP_MAP_FORCE_P): New.
+ (GOMP_MAP_PRESENT_P): New.
+
2022-06-21 Kwok Cheung Yeung <kcy@codesourcery.com>
* cuda/cuda.h (CUdevice_attribute): Add definitions for
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 99684c927c7..7d7bfb425f4 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -42,6 +42,7 @@
#define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4)
#define GOMP_MAP_FLAG_SPECIAL_3 (1 << 5)
#define GOMP_MAP_FLAG_SPECIAL_4 (1 << 6)
+#define GOMP_MAP_FLAG_SPECIAL_5 (1 << 7)
#define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \
| GOMP_MAP_FLAG_SPECIAL_0)
#define GOMP_MAP_DEEP_COPY (GOMP_MAP_FLAG_SPECIAL_4 \
@@ -55,9 +56,14 @@
| GOMP_MAP_FLAG_SPECIAL_1 \
| GOMP_MAP_FLAG_SPECIAL_2 \
| GOMP_MAP_FLAG_SPECIAL_3 \
- | GOMP_MAP_FLAG_SPECIAL_4)
+ | GOMP_MAP_FLAG_SPECIAL_4 \
+ | GOMP_MAP_FLAG_SPECIAL_5)
/* Flag to force a specific behavior (or else, trigger a run-time error). */
-#define GOMP_MAP_FLAG_FORCE (1 << 7)
+#define GOMP_MAP_FLAG_FORCE (GOMP_MAP_FLAG_SPECIAL_5)
+#define GOMP_MAP_FLAG_PRESENT (GOMP_MAP_FLAG_SPECIAL_5 \
+ | GOMP_MAP_FLAG_SPECIAL_0)
+#define GOMP_MAP_FLAG_ALWAYS_PRESENT (GOMP_MAP_FLAG_SPECIAL_2 \
+ | GOMP_MAP_FLAG_PRESENT)
enum gomp_map_kind
{
@@ -130,6 +136,23 @@ 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),
+ /* Must already be present, unconditionally copy from device. */
+ GOMP_MAP_ALWAYS_PRESENT_FROM = (GOMP_MAP_FLAG_ALWAYS_PRESENT
+ | GOMP_MAP_FROM),
+ /* Must already be present, unconditionally copy to and from device. */
+ GOMP_MAP_ALWAYS_PRESENT_TOFROM = (GOMP_MAP_FLAG_ALWAYS_PRESENT
+ | GOMP_MAP_TOFROM),
/* Map a sparse struct; the address is the base of the structure, alignment
it's required alignment, and size is the number of adjacent entries
that belong to the struct. The adjacent entries should be sorted by
@@ -211,11 +234,11 @@ enum gomp_map_kind
};
#define GOMP_MAP_COPY_TO_P(X) \
- (!((X) & GOMP_MAP_FLAG_SPECIAL) \
+ ((!((X) & GOMP_MAP_FLAG_SPECIAL) || GOMP_MAP_PRESENT_P (X)) \
&& ((X) & GOMP_MAP_FLAG_TO))
#define GOMP_MAP_COPY_FROM_P(X) \
- (!((X) & GOMP_MAP_FLAG_SPECIAL) \
+ ((!((X) & GOMP_MAP_FLAG_SPECIAL) || GOMP_MAP_PRESENT_P (X)) \
&& ((X) & GOMP_MAP_FLAG_FROM))
#define GOMP_MAP_ALWAYS_POINTER_P(X) \
@@ -226,17 +249,27 @@ enum gomp_map_kind
|| (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
#define GOMP_MAP_ALWAYS_TO_P(X) \
- (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
+ (((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM) \
+ || ((X) == GOMP_MAP_ALWAYS_PRESENT_TO) \
+ || ((X) == GOMP_MAP_ALWAYS_PRESENT_TOFROM))
#define GOMP_MAP_ALWAYS_FROM_P(X) \
- (((X) == GOMP_MAP_ALWAYS_FROM) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
+ (((X) == GOMP_MAP_ALWAYS_FROM) || ((X) == GOMP_MAP_ALWAYS_TOFROM) \
+ || ((X) == GOMP_MAP_ALWAYS_PRESENT_FROM) \
+ || ((X) == GOMP_MAP_ALWAYS_PRESENT_TOFROM))
#define GOMP_MAP_ALWAYS_P(X) \
- (GOMP_MAP_ALWAYS_TO_P (X) || ((X) == GOMP_MAP_ALWAYS_FROM))
+ (GOMP_MAP_ALWAYS_TO_P (X) || GOMP_MAP_ALWAYS_FROM_P (X))
#define GOMP_MAP_IMPLICIT_P(X) \
(((X) & GOMP_MAP_FLAG_SPECIAL_BITS) == GOMP_MAP_IMPLICIT)
+#define GOMP_MAP_FORCE_P(X) \
+ (((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)
+
#define GOMP_MAP_NONCONTIG_ARRAY_P(X) \
((X) & GOMP_MAP_NONCONTIG_ARRAY)
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 4c1a86e4abe..c05d33f39c7 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,19 @@
+2023-02-01 Kwok Cheung Yeung <kcy@codesourcery.com>
+
+ * target.c (gomp_to_device_kind_p): Add map kinds with 'present'
+ modifier.
+ (gomp_map_vars_existing): Use new GOMP_MAP_FORCE_P macro.
+ (gomp_map_vars_internal): Emit runtime error if memory region not
+ present.
+ (gomp_update): Likewise.
+ (gomp_target_rev): Likewise.
+ * testsuite/libgomp.c-c++-common/target-present-1.c: New.
+ * testsuite/libgomp.c-c++-common/target-present-2.c: New.
+ * testsuite/libgomp.c-c++-common/target-present-3.c: New.
+ * testsuite/libgomp.fortran/target-present-1.f90: New.
+ * testsuite/libgomp.fortran/target-present-2.f90: New.
+ * testsuite/libgomp.fortran/target-present-3.f90: New.
+
2023-02-09 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.fortran/allocate-5.f90: Paste.
diff --git a/libgomp/target.c b/libgomp/target.c
index bbe1a7c4899..b9db0d18e44 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -358,6 +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:
return false;
default:
return true;
@@ -593,7 +595,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
else
tgt_var->length = newn->host_end - newn->host_start;
- if ((kind & GOMP_MAP_FLAG_FORCE)
+ if (GOMP_MAP_FORCE_P (kind)
/* For implicit maps, old contained in new is valid. */
|| !(implicit_subset
/* Otherwise, new contained inside old is considered valid. */
@@ -1805,6 +1807,20 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
#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: !omp_target_is_present "
+ "(%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,
@@ -2321,6 +2337,20 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size);
}
}
+ else
+ {
+ int kind = get_kind (short_mapkind, kinds, i);
+
+ if (GOMP_MAP_PRESENT_P (kind))
+ {
+ /* We already looked up the memory region above and it
+ was missing. */
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("present clause: !omp_target_is_present "
+ "(%p, %d)",
+ (void *) hostaddrs[i], devicep->target_id);
+ }
+ }
}
gomp_mutex_unlock (&devicep->lock);
}
@@ -3629,7 +3659,8 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
case GOMP_MAP_DELETE:
case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION:
- /* Assume it is present; look it up - but ignore otherwise. */
+ /* Assume it is present; look it up - but ignore unless the
+ present clause is there. */
case GOMP_MAP_ALLOC:
case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_ALLOC:
@@ -3641,6 +3672,12 @@ 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_ALWAYS_PRESENT_FROM:
+ case GOMP_MAP_ALWAYS_PRESENT_TO:
+ case GOMP_MAP_ALWAYS_PRESENT_TOFROM:
case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
cdata[i].devaddr = devaddrs[i];
bool zero_len = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
@@ -3661,7 +3698,23 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
devaddrs[i] + sizes[i], zero_len);
cdata[i].present = n2 != NULL;
}
- if (!cdata[i].present
+ if (!cdata[i].present && GOMP_MAP_PRESENT_P (kind))
+ {
+ gomp_mutex_unlock (&devicep->lock);
+#ifdef HAVE_INTTYPES_H
+ gomp_fatal ("present clause: no corresponding data on "
+ "parent device at %p with size %"PRIu64,
+ (void *) (uintptr_t) devaddrs[i],
+ (uint64_t) sizes[i]);
+#else
+ gomp_fatal ("present clause: no corresponding data on "
+ "parent device at %p with size %lu",
+ (void *) (uintptr_t) devaddrs[i],
+ (unsigned long) sizes[i]);
+#endif
+ break;
+ }
+ else if (!cdata[i].present
&& kind != GOMP_MAP_DELETE
&& kind != GOMP_MAP_RELEASE
&& kind != GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION)
@@ -3679,8 +3732,7 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr,
&& (kind == GOMP_MAP_TO || kind == GOMP_MAP_TOFROM))
|| kind == GOMP_MAP_FORCE_TO
|| kind == GOMP_MAP_FORCE_TOFROM
- || kind == GOMP_MAP_ALWAYS_TO
- || kind == GOMP_MAP_ALWAYS_TOFROM)
+ || GOMP_MAP_ALWAYS_TO_P (kind))
{
if (dev_to_host_cpy)
dev_to_host_cpy ((void *) (uintptr_t) devaddrs[i],
@@ -3865,6 +3917,10 @@ 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;
/* FALLTHRU */
case GOMP_MAP_FROM:
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
new file mode 100644
index 00000000000..bbc4559b12e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c
@@ -0,0 +1,27 @@
+/* { dg-do run { target offload_target_any } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#define N 100
+
+int main (void)
+{
+ int a[N], b[N], c[N];
+
+ for (int i = 0; i < N; i++) {
+ a[i] = i * 2;
+ b[i] = i * 3 + 1;
+ }
+
+ #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)
+ for (int i = 0; i < N; i++)
+ c[i] = a[i];
+
+ /* b has not been allocated, so this should result in an error. */
+ /* { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" } */
+ #pragma omp target map (present, to: b)
+ for (int i = 0; i < N; i++)
+ c[i] += b[i];
+ #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
new file mode 100644
index 00000000000..6259c959c04
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c
@@ -0,0 +1,27 @@
+/* { dg-do run { target offload_target_any } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#define N 100
+
+int main (void)
+{
+ int a[N], b[N], c[N];
+
+ for (int i = 0; i < N; i++) {
+ a[i] = i * 2;
+ b[i] = i * 3 + 1;
+ }
+
+ #pragma omp target enter data map (alloc: a, c)
+ /* a has already been allocated, so this should be okay. */
+ #pragma omp target defaultmap (present)
+ for (int i = 0; i < N; i++)
+ c[i] = a[i];
+
+ /* b has not been allocated, so this should result in an error. */
+ /* { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" } */
+ #pragma omp target defaultmap (present)
+ for (int i = 0; i < N; i++)
+ c[i] += b[i];
+ #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
new file mode 100644
index 00000000000..89e648645b2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c
@@ -0,0 +1,27 @@
+/* { dg-do run { target offload_target_any } } */
+/* { dg-shouldfail "present error triggered" } */
+
+#include <stdio.h>
+
+#define N 100
+
+int main (void)
+{
+ int a[N], b[N], c[N];
+
+ for (int i = 0; i < N; i++) {
+ a[i] = i * 2;
+ b[i] = i * 3 + 1;
+ }
+
+ #pragma omp target enter data map (alloc: a, c)
+
+ /* This should work as a has already been allocated. */
+ #pragma omp target update to (present: a)
+
+ /* This should fail as b has not been allocated. */
+ /* { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" } */
+ #pragma omp target update to (present: b)
+
+ #pragma omp target exit data map (from: c)
+}
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-1.f90 b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
new file mode 100644
index 00000000000..80046011b25
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-1.f90
@@ -0,0 +1,30 @@
+! { dg-do run { target offload_target_any } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+ implicit none
+ integer, parameter :: N = 100
+ integer :: a(N), b(N), c(N), i
+
+ do i = 1, N
+ a(i) = i * 2
+ b(i) = i * 3 + 1
+ end do
+
+ !$omp target enter data map (alloc: a)
+ ! a has already been allocated, so this should be okay.
+ !$omp target map (present, to: a)
+ do i = 1, N
+ c(i) = a(i)
+ end do
+ !$omp end target
+
+ ! b has not been allocated, so this should result in an error.
+ ! { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" }
+ !$omp target map (present, to: b)
+ do i = 1, N
+ c(i) = c(i) + b(i)
+ end do
+ !$omp end target
+ !$omp target exit data map (from: c)
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-2.f90 b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
new file mode 100644
index 00000000000..0a38dea1e41
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-2.f90
@@ -0,0 +1,30 @@
+! { dg-do run { target offload_target_any } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+ implicit none
+ integer, parameter :: N = 100
+ integer :: a(N), b(N), c(N), i
+
+ do i = 1, N
+ a(i) = i * 2
+ b(i) = i * 3 + 1
+ end do
+
+ !$omp target enter data map (alloc: a)
+ ! a has already been allocated, so this should be okay.
+ !$omp target defaultmap (present)
+ do i = 1, N
+ c(i) = a(i)
+ end do
+ !$omp end target
+
+ ! b has not been allocated, so this should result in an error.
+ ! { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" }
+ !$omp target defaultmap (present)
+ do i = 1, N
+ c(i) = c(i) + b(i)
+ end do
+ !$omp end target
+!$omp target exit data map (from: c)
+end program
diff --git a/libgomp/testsuite/libgomp.fortran/target-present-3.f90 b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
new file mode 100644
index 00000000000..c4deb8652d1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/target-present-3.f90
@@ -0,0 +1,22 @@
+! { dg-do run { target offload_target_any } }
+! { dg-shouldfail "present error triggered" }
+
+program main
+ implicit none
+ integer, parameter :: N = 100
+ integer :: a(N), b(N), c(N), i
+
+ do i = 1, N
+ a(i) = i * 2
+ b(i) = i * 3 + 1
+ 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 fail as b has not been allocated.
+ ! { dg-output "libgomp: present clause: !omp_target_is_present \\\(0x\[0-9a-f\]+, \[0-9\]+\\\)" }
+ !$omp target update to (present: b)
+ !$omp target exit data map (from: c)
+end program
^ permalink raw reply [flat|nested] 2+ messages in thread
end of thread, other threads:[~2023-06-09 11:18 UTC | newest]
Thread overview: 2+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-06-09 11:18 [gcc/devel/omp/gcc-13] openmp: Add support for the 'present' modifier Tobias Burnus
-- strict thread matches above, loose matches on Subject: below --
2023-05-19 16:52 Kwok Yeung
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).