diff options
Diffstat (limited to 'gcc')
23 files changed, 797 insertions, 51 deletions
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 5baa501..72f6fba 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -14939,6 +14939,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; @@ -17109,6 +17116,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); @@ -17140,11 +17148,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; } @@ -17156,14 +17174,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) @@ -17419,21 +17448,42 @@ 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, 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, 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); + 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: @@ -17940,11 +17990,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: @@ -21768,11 +21820,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: @@ -21918,7 +21977,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: @@ -21929,6 +21991,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: @@ -22016,6 +22086,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; @@ -22028,6 +22100,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: @@ -22273,11 +22353,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/parser.cc b/gcc/cp/parser.cc index 1c9aa671..d77fbd2 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -38779,6 +38779,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; @@ -40487,6 +40494,41 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc) } /* 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, 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 ) @@ -40532,6 +40574,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); @@ -40568,11 +40611,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, @@ -40588,15 +40644,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 @@ -41373,13 +41439,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, - 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, - true); + clauses = cp_parser_omp_clause_from_to (parser, OMP_CLAUSE_FROM, + clauses); c_name = "from"; break; case PRAGMA_OMP_CLAUSE_UNIFORM: @@ -45231,11 +45297,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: @@ -45338,7 +45411,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: @@ -45349,6 +45425,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: @@ -45441,6 +45525,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; @@ -45453,6 +45539,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 0843b5e..c045146 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -10030,11 +10030,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/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 6d75cc2..99c8bda 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1468,9 +1468,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 3e5f942..33ca498 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1308,7 +1308,14 @@ enum gfc_omp_map_op OMP_MAP_RELEASE, OMP_MAP_ALWAYS_TO, OMP_MAP_ALWAYS_FROM, - OMP_MAP_ALWAYS_TOFROM + 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 }; enum gfc_omp_defaultmap @@ -1362,6 +1369,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 4c30548..8efc4b3 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -1336,6 +1336,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 ) */ @@ -2098,6 +2122,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) @@ -2105,7 +2131,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 ()) @@ -2529,10 +2555,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; break; case 'g': @@ -2888,8 +2912,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 (;;) { @@ -2904,20 +2930,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) @@ -2941,6 +2985,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], @@ -3467,10 +3517,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': @@ -8092,11 +8140,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, " @@ -8112,7 +8167,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; @@ -8120,6 +8178,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 " @@ -8132,6 +8196,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; @@ -8141,6 +8207,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 42b608f..4aa16fa 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -3099,6 +3099,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; @@ -3894,6 +3918,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; @@ -4435,6 +4461,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 4aa6229..91640de 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -7926,6 +7926,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); @@ -9100,6 +9105,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: @@ -9331,6 +9343,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: @@ -10812,6 +10831,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) { @@ -11968,6 +12031,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) { @@ -12412,6 +12478,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; default: gcc_unreachable (); } diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 9db33d2..1857b5b 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1576,6 +1576,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 @@ -12797,6 +12800,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: @@ -13338,6 +13349,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: @@ -13377,11 +13395,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/c-c++-common/gomp/defaultmap-4.c b/gcc/testsuite/c-c++-common/gomp/defaultmap-4.c new file mode 100644 index 0000000..1afff7e --- /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 6ee5971..8c5b7f7 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 0000000..4b4bd6d --- /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 0000000..4408ab7 --- /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 299d971..1f1b852 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 0000000..669a623 --- /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 0000000..9eb956f --- /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 0000000..74bd01f --- /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 009c6d4..317090a 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 92b802c..15ebdd6 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 0000000..f99bffe --- /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\\(present:b \\\[len: \[0-9\]+\\\]\\) from\\(d \\\[len: \[0-9\]+\\\]\\)" "gimple" } } diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 6dd7b68..c48a12b 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -572,7 +572,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 db2a58a..25d191b 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -991,6 +991,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; default: gcc_unreachable (); } @@ -1038,12 +1059,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; @@ -1210,6 +1235,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; @@ -1773,6 +1773,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. */ |