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;
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);
}
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;
}
&& 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)
}
/* 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:
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:
{
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:
{
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:
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:
{
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;
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:
{
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:
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;
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, 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 )
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);
}
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,
{
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
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:
{
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:
{
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:
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:
{
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;
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:
{
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:
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;
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);
}
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
} linear;
struct gfc_common_head *common;
bool lastprivate_conditional;
+ bool present_modifier;
} u;
union
{
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 ) */
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)
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 ())
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':
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 (;;)
{
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)
&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],
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':
{
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, "
{
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;
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 "
{
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;
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 "
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;
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;
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;
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);
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:
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:
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)
{
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)
{
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 ();
}
&& 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
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:
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:
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:
--- /dev/null
+/* { 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" } } */
#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 } } */
#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)
;
/* { 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" } } */
--- /dev/null
+/* { 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" } } */
--- /dev/null
+/* { 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" } } */
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" }
--- /dev/null
+! { 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" } }
--- /dev/null
+! { 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" } }
--- /dev/null
+! { 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" } }
implicit none
-integer :: a, b, close, always, to
+integer :: a, b, close, always, to, present
!$omp target map(close)
!$omp end target
!$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" } }
!$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
--- /dev/null
+! { 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" } }
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 {
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 ();
}
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;
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;
(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. */
#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 \
| 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
{
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
};
#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) \
|| (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)
+
/* Asynchronous behavior. Keep in sync with
libgomp/{openacc.h,openacc.f90,openacc_lib.h}:acc_async_t. */
@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
@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
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;
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. */
#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,
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);
}
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:
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
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)
&& (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],
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:
--- /dev/null
+/* { 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)
+}
--- /dev/null
+/* { 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)
+}
--- /dev/null
+/* { 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)
+}
--- /dev/null
+! { 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
--- /dev/null
+! { 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
--- /dev/null
+! { 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
projects / thirdparty / gcc.git / commitdiff
