mirror of git://gcc.gnu.org/git/gcc.git
openmp: Add support for iterators in 'target update' clauses (C/C++)
This adds support for iterators in 'to' and 'from' clauses in the 'target update' OpenMP directive. gcc/c/ * c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator' modifier. * c-typeck.cc (c_finish_omp_clauses): Finish iterators for to/from clauses. gcc/cp/ * parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator' modifier. * semantics.cc (finish_omp_clauses): Finish iterators for to/from clauses. gcc/ * gimplify.cc (remove_unused_omp_iterator_vars): Display unused variable warning for 'to' and 'from' clauses. (gimplify_scan_omp_clauses): Add argument for iterator loop sequence. Gimplify the clause decl and size into the iterator loop if iterators are used. (gimplify_omp_workshare): Add argument for iterator loops sequence in call to gimplify_scan_omp_clauses. (gimplify_omp_target_update): Call remove_unused_omp_iterator_vars and build_omp_iterators_loops. Add loop sequence as argument when calling gimplify_scan_omp_clauses, gimplify_adjust_omp_clauses and building the Gimple statement. * tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators for to/from clauses with iterators. * tree.cc (omp_clause_num_ops): Add extra operand for OMP_CLAUSE_FROM and OMP_CLAUSE_TO. * tree.h (OMP_CLAUSE_HAS_ITERATORS): Add check for OMP_CLAUSE_TO and OMP_CLAUSE_FROM. (OMP_CLAUSE_ITERATORS): Likewise. gcc/testsuite/ * c-c++-common/gomp/target-update-iterators-1.c: New. * c-c++-common/gomp/target-update-iterators-2.c: New. * c-c++-common/gomp/target-update-iterators-3.c: New. libgomp/ * target.c (gomp_update): Call gomp_merge_iterator_maps. Free allocated variables. * testsuite/libgomp.c-c++-common/target-update-iterators-1.c: New. * testsuite/libgomp.c-c++-common/target-update-iterators-2.c: New. * testsuite/libgomp.c-c++-common/target-update-iterators-3.c: New.
This commit is contained in:
Kwok Cheung Yeung
parent
8b8b0eada6
commit
87262627fd
|
|
@ -20576,8 +20576,11 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list)
|
|||
to ( variable-list )
|
||||
|
||||
OpenMP 5.1:
|
||||
from ( [present :] variable-list )
|
||||
to ( [present :] variable-list ) */
|
||||
from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
|
||||
to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
|
||||
|
||||
motion-modifier:
|
||||
present | iterator (iterators-definition) */
|
||||
|
||||
static tree
|
||||
c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
|
||||
|
|
@ -20588,18 +20591,85 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
|
|||
if (!parens.require_open (parser))
|
||||
return list;
|
||||
|
||||
bool present = false;
|
||||
c_token *token = c_parser_peek_token (parser);
|
||||
int pos = 1, colon_pos = 0;
|
||||
int iterator_length = 0;
|
||||
|
||||
if (token->type == CPP_NAME
|
||||
&& strcmp (IDENTIFIER_POINTER (token->value), "present") == 0
|
||||
&& c_parser_peek_2nd_token (parser)->type == CPP_COLON)
|
||||
while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
|
||||
{
|
||||
present = true;
|
||||
c_parser_consume_token (parser);
|
||||
c_parser_consume_token (parser);
|
||||
const char *identifier =
|
||||
IDENTIFIER_POINTER (c_parser_peek_nth_token_raw (parser, pos)->value);
|
||||
if (c_parser_peek_nth_token_raw (parser, pos + 1)->type
|
||||
== CPP_OPEN_PAREN)
|
||||
{
|
||||
unsigned int npos = pos + 2;
|
||||
if (c_parser_check_balanced_raw_token_sequence (parser, &npos)
|
||||
&& (c_parser_peek_nth_token_raw (parser, npos)->type
|
||||
== CPP_CLOSE_PAREN))
|
||||
{
|
||||
if (strcmp (identifier, "iterator") == 0)
|
||||
iterator_length = npos - pos + 1;
|
||||
pos = npos;
|
||||
}
|
||||
}
|
||||
if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
|
||||
pos += 2;
|
||||
else
|
||||
pos++;
|
||||
if (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COLON)
|
||||
{
|
||||
colon_pos = pos;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
bool present = false;
|
||||
tree iterators = NULL_TREE;
|
||||
|
||||
for (int pos = 1; pos < colon_pos; ++pos)
|
||||
{
|
||||
c_token *token = c_parser_peek_token (parser);
|
||||
if (token->type == CPP_COMMA)
|
||||
{
|
||||
c_parser_consume_token (parser);
|
||||
continue;
|
||||
}
|
||||
const char *p = IDENTIFIER_POINTER (token->value);
|
||||
if (strcmp ("present", p) == 0)
|
||||
{
|
||||
if (present)
|
||||
{
|
||||
c_parser_error (parser, "too many %<present%> modifiers");
|
||||
parens.skip_until_found_close (parser);
|
||||
return list;
|
||||
}
|
||||
present = true;
|
||||
c_parser_consume_token (parser);
|
||||
}
|
||||
else if (strcmp ("iterator", p) == 0)
|
||||
{
|
||||
if (iterators)
|
||||
{
|
||||
c_parser_error (parser, "too many %<iterator%> modifiers");
|
||||
parens.skip_until_found_close (parser);
|
||||
return list;
|
||||
}
|
||||
iterators = c_parser_omp_iterators (parser);
|
||||
pos += iterator_length - 1;
|
||||
}
|
||||
else
|
||||
{
|
||||
error_at (token->location,
|
||||
"%qs clause with modifier other than %<iterator%> or "
|
||||
"%<present%>",
|
||||
kind == OMP_CLAUSE_TO ? "to" : "from");
|
||||
parens.skip_until_found_close (parser);
|
||||
return list;
|
||||
}
|
||||
}
|
||||
|
||||
if (colon_pos)
|
||||
c_parser_require (parser, CPP_COLON, "expected %<:%>");
|
||||
|
||||
tree nl = c_parser_omp_variable_list (parser, loc, kind, list);
|
||||
parens.skip_until_found_close (parser);
|
||||
|
||||
|
|
@ -20607,6 +20677,19 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
|
|||
for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
|
||||
OMP_CLAUSE_MOTION_PRESENT (c) = 1;
|
||||
|
||||
if (iterators)
|
||||
{
|
||||
tree block = pop_scope ();
|
||||
if (iterators == error_mark_node)
|
||||
iterators = NULL_TREE;
|
||||
else
|
||||
TREE_VEC_ELT (iterators, 5) = block;
|
||||
}
|
||||
|
||||
if (iterators)
|
||||
for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
|
||||
OMP_CLAUSE_ITERATORS (c) = iterators;
|
||||
|
||||
return nl;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -16966,6 +16966,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
|||
remove = true;
|
||||
break;
|
||||
}
|
||||
/* FALLTHRU */
|
||||
case OMP_CLAUSE_TO:
|
||||
case OMP_CLAUSE_FROM:
|
||||
if (OMP_CLAUSE_ITERATORS (c)
|
||||
&& c_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c)))
|
||||
{
|
||||
|
|
@ -16973,8 +16976,6 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
|||
break;
|
||||
}
|
||||
/* FALLTHRU */
|
||||
case OMP_CLAUSE_TO:
|
||||
case OMP_CLAUSE_FROM:
|
||||
case OMP_CLAUSE__CACHE_:
|
||||
{
|
||||
using namespace omp_addr_tokenizer;
|
||||
|
|
|
|||
113
gcc/cp/parser.cc
113
gcc/cp/parser.cc
|
|
@ -42636,8 +42636,11 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc)
|
|||
to ( variable-list )
|
||||
|
||||
OpenMP 5.1:
|
||||
from ( [present :] variable-list )
|
||||
to ( [present :] variable-list ) */
|
||||
from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
|
||||
to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
|
||||
|
||||
motion-modifier:
|
||||
present | iterator (iterators-definition) */
|
||||
|
||||
static tree
|
||||
cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
|
||||
|
|
@ -42646,23 +42649,113 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
|
|||
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);
|
||||
int pos = 1;
|
||||
int colon_pos = 0;
|
||||
int iterator_length = 0;
|
||||
|
||||
if (token->type == CPP_NAME
|
||||
&& strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0
|
||||
&& cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
|
||||
while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME)
|
||||
{
|
||||
present = true;
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
const char *identifier =
|
||||
IDENTIFIER_POINTER (cp_lexer_peek_nth_token (parser->lexer,
|
||||
pos)->u.value);
|
||||
if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_OPEN_PAREN))
|
||||
{
|
||||
int n = cp_parser_skip_balanced_tokens (parser, pos + 1);
|
||||
if (n != pos + 1)
|
||||
{
|
||||
if (strcmp (identifier, "iterator") == 0)
|
||||
iterator_length = n - pos;
|
||||
pos = n - 1;
|
||||
}
|
||||
}
|
||||
if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
|
||||
pos += 2;
|
||||
else
|
||||
pos++;
|
||||
if (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_COLON)
|
||||
{
|
||||
colon_pos = pos;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
bool present = false;
|
||||
tree iterators = NULL_TREE;
|
||||
|
||||
for (int pos = 1; pos < colon_pos; ++pos)
|
||||
{
|
||||
cp_token *token = cp_lexer_peek_token (parser->lexer);
|
||||
if (token->type == CPP_COMMA)
|
||||
{
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
continue;
|
||||
}
|
||||
const char *p = IDENTIFIER_POINTER (token->u.value);
|
||||
if (strcmp ("present", p) == 0)
|
||||
{
|
||||
if (present)
|
||||
{
|
||||
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 = true;
|
||||
cp_lexer_consume_token (parser->lexer);
|
||||
}
|
||||
else if (strcmp ("iterator", p) == 0)
|
||||
{
|
||||
if (iterators)
|
||||
{
|
||||
cp_parser_error (parser, "too many %<iterator%> modifiers");
|
||||
cp_parser_skip_to_closing_parenthesis (parser,
|
||||
/*recovering=*/true,
|
||||
/*or_comma=*/false,
|
||||
/*consume_paren=*/true);
|
||||
return list;
|
||||
}
|
||||
begin_scope (sk_omp, NULL);
|
||||
iterators = cp_parser_omp_iterators (parser);
|
||||
pos += iterator_length - 1;
|
||||
}
|
||||
|
||||
else
|
||||
{
|
||||
error_at (token->location,
|
||||
"%qs clause with modifier other than %<iterator%> "
|
||||
"or %<present%>",
|
||||
kind == OMP_CLAUSE_TO ? "to" : "from");
|
||||
cp_parser_skip_to_closing_parenthesis (parser,
|
||||
/*recovering=*/true,
|
||||
/*or_comma=*/false,
|
||||
/*consume_paren=*/true);
|
||||
return list;
|
||||
}
|
||||
}
|
||||
|
||||
if (colon_pos)
|
||||
cp_parser_require (parser, CPP_COLON, RT_COLON);
|
||||
|
||||
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;
|
||||
|
||||
if (iterators)
|
||||
{
|
||||
tree block = poplevel (1, 1, 0);
|
||||
if (iterators == error_mark_node)
|
||||
iterators = NULL_TREE;
|
||||
else
|
||||
TREE_VEC_ELT (iterators, 5) = block;
|
||||
}
|
||||
|
||||
if (iterators)
|
||||
for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
|
||||
OMP_CLAUSE_ITERATORS (c) = iterators;
|
||||
|
||||
return nl;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -9007,6 +9007,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
|||
remove = true;
|
||||
break;
|
||||
}
|
||||
/* FALLTHRU */
|
||||
case OMP_CLAUSE_TO:
|
||||
case OMP_CLAUSE_FROM:
|
||||
if (OMP_CLAUSE_ITERATORS (c)
|
||||
&& cp_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c)))
|
||||
{
|
||||
|
|
@ -9014,8 +9017,6 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
|||
break;
|
||||
}
|
||||
/* FALLTHRU */
|
||||
case OMP_CLAUSE_TO:
|
||||
case OMP_CLAUSE_FROM:
|
||||
case OMP_CLAUSE__CACHE_:
|
||||
{
|
||||
using namespace omp_addr_tokenizer;
|
||||
|
|
|
|||
|
|
@ -9966,9 +9966,11 @@ remove_unused_omp_iterator_vars (tree *list_p)
|
|||
if (t == NULL_TREE)
|
||||
{
|
||||
need_new_iterators = true;
|
||||
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|
||||
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO
|
||||
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM))
|
||||
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|
||||
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO
|
||||
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM))
|
||||
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|
||||
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM)
|
||||
warning_at (OMP_CLAUSE_LOCATION (c), OPT_Wopenmp,
|
||||
"iterator variable %qE not used in clause "
|
||||
"expression", DECL_NAME (var));
|
||||
|
|
@ -13582,7 +13584,8 @@ omp_instantiate_implicit_mappers (splay_tree_node n, void *data)
|
|||
static void
|
||||
gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|
||||
enum omp_region_type region_type,
|
||||
enum tree_code code)
|
||||
enum tree_code code,
|
||||
gimple_seq *loops_seq_p = NULL)
|
||||
{
|
||||
using namespace omp_addr_tokenizer;
|
||||
struct gimplify_omp_ctx *ctx, *outer_ctx;
|
||||
|
|
@ -14353,23 +14356,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|
|||
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
|
||||
OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
|
||||
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
|
||||
if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
|
||||
NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
|
||||
gimple_seq *seq_p;
|
||||
seq_p = enter_omp_iterator_loop_context (c, loops_seq_p, pre_p);
|
||||
if (gimplify_expr (&OMP_CLAUSE_SIZE (c), seq_p, NULL,
|
||||
is_gimple_val, fb_rvalue) == GS_ERROR)
|
||||
{
|
||||
remove = true;
|
||||
exit_omp_iterator_loop_context (c);
|
||||
break;
|
||||
}
|
||||
if (!DECL_P (decl))
|
||||
{
|
||||
if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
|
||||
NULL, is_gimple_lvalue, fb_lvalue)
|
||||
== GS_ERROR)
|
||||
{
|
||||
remove = true;
|
||||
break;
|
||||
}
|
||||
if (gimplify_expr (&OMP_CLAUSE_DECL (c), seq_p, NULL,
|
||||
is_gimple_lvalue, fb_lvalue) == GS_ERROR)
|
||||
remove = true;
|
||||
exit_omp_iterator_loop_context (c);
|
||||
break;
|
||||
}
|
||||
exit_omp_iterator_loop_context (c);
|
||||
goto do_notice;
|
||||
|
||||
case OMP_CLAUSE__MAPPER_BINDING_:
|
||||
|
|
@ -18728,7 +18732,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
|
|||
if ((ort & ORT_ACC) == 0)
|
||||
in_omp_construct = false;
|
||||
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort,
|
||||
TREE_CODE (expr));
|
||||
TREE_CODE (expr), &iterator_loops_seq);
|
||||
if (TREE_CODE (expr) == OMP_TARGET)
|
||||
optimize_target_teams (expr, pre_p);
|
||||
if ((ort & (ORT_TARGET | ORT_TARGET_DATA)) != 0
|
||||
|
|
@ -18885,10 +18889,16 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
|
|||
default:
|
||||
gcc_unreachable ();
|
||||
}
|
||||
|
||||
gimple_seq iterator_loops_seq = NULL;
|
||||
remove_unused_omp_iterator_vars (&OMP_STANDALONE_CLAUSES (expr));
|
||||
build_omp_iterators_loops (&OMP_STANDALONE_CLAUSES (expr),
|
||||
&iterator_loops_seq);
|
||||
|
||||
gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p,
|
||||
ort, TREE_CODE (expr));
|
||||
ort, TREE_CODE (expr), &iterator_loops_seq);
|
||||
gimplify_adjust_omp_clauses (pre_p, NULL, &OMP_STANDALONE_CLAUSES (expr),
|
||||
TREE_CODE (expr));
|
||||
TREE_CODE (expr), &iterator_loops_seq);
|
||||
if (TREE_CODE (expr) == OACC_UPDATE
|
||||
&& omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
|
||||
OMP_CLAUSE_IF_PRESENT))
|
||||
|
|
@ -18952,7 +18962,8 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
|
|||
gcc_unreachable ();
|
||||
}
|
||||
}
|
||||
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
|
||||
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr),
|
||||
iterator_loops_seq);
|
||||
|
||||
gimplify_seq_add_stmt (pre_p, stmt);
|
||||
*expr_p = NULL_TREE;
|
||||
|
|
|
|||
|
|
@ -0,0 +1,21 @@
|
|||
/* { dg-do compile } */
|
||||
/* { dg-options "-fopenmp" } */
|
||||
|
||||
#define DIM1 17
|
||||
#define DIM2 39
|
||||
|
||||
void f (int **x, float **y)
|
||||
{
|
||||
#pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2])
|
||||
|
||||
#pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2], y[i][:DIM2])
|
||||
|
||||
#pragma omp target update to (iterator(i=0:DIM1), present: x[i][:DIM2])
|
||||
|
||||
#pragma omp target update to (iterator(i=0:DIM1), iterator(j=0:DIM2): x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */
|
||||
/* { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-1 } */
|
||||
|
||||
#pragma omp target update from (iterator(i=0:DIM1), something: x[i][j]) /* { dg-error "'from' clause with modifier other than 'iterator' or 'present'" } */
|
||||
/* { dg-error "expected '\\)' before 'something'" "" { target c } .-1 } */
|
||||
/* { dg-error "'#pragma omp target update' must contain at least one 'from' or 'to' clauses" "" { target *-*-* } .-2 } */
|
||||
}
|
||||
|
|
@ -0,0 +1,23 @@
|
|||
/* { dg-do compile } */
|
||||
/* { dg-options "-fopenmp -fdump-tree-gimple" } */
|
||||
|
||||
void f (int *x, float *y, double *z)
|
||||
{
|
||||
#pragma omp target update to(iterator(i=0:10): x) /* { dg-warning "iterator variable 'i' not used in clause expression" }*/
|
||||
;
|
||||
|
||||
#pragma omp target update from(iterator(i2=0:10, j2=0:20): x[i2]) /* { dg-warning "iterator variable 'j2' not used in clause expression" }*/
|
||||
;
|
||||
|
||||
#pragma omp target update to(iterator(i3=0:10, j3=0:20, k3=0:30): x[i3+j3], y[j3+k3], z[k3+i3])
|
||||
/* { dg-warning "iterator variable 'i3' not used in clause expression" "" { target *-*-* } .-1 } */
|
||||
/* { dg-warning "iterator variable 'j3' not used in clause expression" "" { target *-*-* } .-2 } */
|
||||
/* { dg-warning "iterator variable 'k3' not used in clause expression" "" { target *-*-* } .-3 } */
|
||||
;
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump "update to\\\(x " "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump "update from\\\(iterator\\\(int i2=0:10:1, loop_label=" "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump "to\\\(iterator\\\(int i3=0:10:1, int k3=0:30:1, loop_label=" "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump "to\\\(iterator\\\(int j3=0:20:1, int k3=0:30:1, loop_label=" "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump "to\\\(iterator\\\(int i3=0:10:1, int j3=0:20:1, loop_label=" "gimple" } } */
|
||||
|
|
@ -0,0 +1,17 @@
|
|||
/* { dg-do compile } */
|
||||
/* { dg-options "-fopenmp -fdump-tree-gimple" } */
|
||||
|
||||
#define DIM1 10
|
||||
#define DIM2 20
|
||||
#define DIM3 30
|
||||
|
||||
void f (int ***x, float ***y, double **z)
|
||||
{
|
||||
#pragma omp target update to (iterator(i=0:DIM1, j=0:DIM2): x[i][j][:DIM3], y[i][j][:DIM3])
|
||||
#pragma omp target update from (iterator(i=0:DIM1): z[i][:DIM2])
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "to\\(iterator\\(int i=0:10:1, int j=0:20:1, loop_label=<D\.\[0-9\]+>, elems=omp_iter_data\.\[0-9\]+, index=D\.\[0-9\]+\\):\\*D\.\[0-9\]+" 2 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1, loop_label=<D\.\[0-9\]+>, elems=omp_iter_data\.\[0-9\]+, index=D\.\[0-9\]+\\):\\*D\.\[0-9\]+" 1 "gimple" } } */
|
||||
|
|
@ -1199,6 +1199,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
|
|||
pp_string (pp, "from(");
|
||||
if (OMP_CLAUSE_MOTION_PRESENT (clause))
|
||||
pp_string (pp, "present:");
|
||||
if (OMP_CLAUSE_ITERATORS (clause))
|
||||
{
|
||||
dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags);
|
||||
pp_colon (pp);
|
||||
}
|
||||
dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
|
||||
spc, flags, false);
|
||||
goto print_clause_size;
|
||||
|
|
@ -1207,6 +1212,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
|
|||
pp_string (pp, "to(");
|
||||
if (OMP_CLAUSE_MOTION_PRESENT (clause))
|
||||
pp_string (pp, "present:");
|
||||
if (OMP_CLAUSE_ITERATORS (clause))
|
||||
{
|
||||
dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags);
|
||||
pp_colon (pp);
|
||||
}
|
||||
dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
|
||||
spc, flags, false);
|
||||
goto print_clause_size;
|
||||
|
|
|
|||
|
|
@ -323,8 +323,8 @@ unsigned const char omp_clause_num_ops[] =
|
|||
1, /* OMP_CLAUSE_IS_DEVICE_PTR */
|
||||
1, /* OMP_CLAUSE_INCLUSIVE */
|
||||
1, /* OMP_CLAUSE_EXCLUSIVE */
|
||||
2, /* OMP_CLAUSE_FROM */
|
||||
2, /* OMP_CLAUSE_TO */
|
||||
3, /* OMP_CLAUSE_FROM */
|
||||
3, /* OMP_CLAUSE_TO */
|
||||
3, /* OMP_CLAUSE_MAP (update walk_tree_1 if this is changed) */
|
||||
1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */
|
||||
1, /* OMP_CLAUSE_DOACROSS */
|
||||
|
|
|
|||
|
|
@ -1660,11 +1660,13 @@ class auto_suppress_location_wrappers
|
|||
#define OMP_CLAUSE_LOCATION(NODE) (OMP_CLAUSE_CHECK (NODE))->omp_clause.locus
|
||||
|
||||
#define OMP_CLAUSE_HAS_ITERATORS(NODE) \
|
||||
(OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_MAP \
|
||||
((OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_FROM \
|
||||
|| OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_TO \
|
||||
|| OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_MAP) \
|
||||
&& OMP_CLAUSE_ITERATORS (NODE))
|
||||
#define OMP_CLAUSE_ITERATORS(NODE) \
|
||||
OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \
|
||||
OMP_CLAUSE_MAP, \
|
||||
OMP_CLAUSE_FROM, \
|
||||
OMP_CLAUSE_MAP), 2)
|
||||
|
||||
/* True on OMP_FOR and other OpenMP/OpenACC looping constructs if the loop nest
|
||||
|
|
|
|||
|
|
@ -2372,6 +2372,8 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
|
|||
size_t i;
|
||||
struct splay_tree_key_s cur_node;
|
||||
const int typemask = short_mapkind ? 0xff : 0x7;
|
||||
bool iterators_p = false;
|
||||
size_t *iterator_count = NULL;
|
||||
|
||||
if (!devicep)
|
||||
return;
|
||||
|
|
@ -2379,6 +2381,10 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
|
|||
if (mapnum == 0)
|
||||
return;
|
||||
|
||||
if (short_mapkind) /* OpenMP */
|
||||
iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
|
||||
&kinds, &iterator_count);
|
||||
|
||||
gomp_mutex_lock (&devicep->lock);
|
||||
if (devicep->state == GOMP_DEVICE_FINALIZED)
|
||||
{
|
||||
|
|
@ -2472,6 +2478,14 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
|
|||
}
|
||||
}
|
||||
gomp_mutex_unlock (&devicep->lock);
|
||||
|
||||
if (iterators_p)
|
||||
{
|
||||
free (hostaddrs);
|
||||
free (sizes);
|
||||
free (kinds);
|
||||
free (iterator_count);
|
||||
}
|
||||
}
|
||||
|
||||
static struct gomp_offload_icv_list *
|
||||
|
|
|
|||
|
|
@ -0,0 +1,65 @@
|
|||
/* { dg-do run } */
|
||||
|
||||
/* Test target enter data and target update to the target using map
|
||||
iterators. */
|
||||
|
||||
#include <stdlib.h>
|
||||
|
||||
#define DIM1 8
|
||||
#define DIM2 15
|
||||
|
||||
int mkarray (int *x[])
|
||||
{
|
||||
int expected = 0;
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
{
|
||||
x[i] = (int *) malloc (DIM2 * sizeof (int));
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
{
|
||||
x[i][j] = rand ();
|
||||
expected += x[i][j];
|
||||
}
|
||||
}
|
||||
|
||||
return expected;
|
||||
}
|
||||
|
||||
int main (void)
|
||||
{
|
||||
int *x[DIM1];
|
||||
int sum;
|
||||
int expected = mkarray (x);
|
||||
|
||||
#pragma omp target enter data map(to: x[:DIM1])
|
||||
#pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
|
||||
#pragma omp target map(from: sum)
|
||||
{
|
||||
sum = 0;
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
sum += x[i][j];
|
||||
}
|
||||
|
||||
if (sum != expected)
|
||||
return 1;
|
||||
|
||||
expected = 0;
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
{
|
||||
x[i][j] *= rand ();
|
||||
expected += x[i][j];
|
||||
}
|
||||
|
||||
#pragma omp target update to(iterator(i=0:DIM1): x[i][:DIM2])
|
||||
|
||||
#pragma omp target map(from: sum)
|
||||
{
|
||||
sum = 0;
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
sum += x[i][j];
|
||||
}
|
||||
|
||||
return sum != expected;
|
||||
}
|
||||
|
|
@ -0,0 +1,58 @@
|
|||
/* { dg-do run } */
|
||||
/* { dg-require-effective-target offload_device_nonshared_as } */
|
||||
|
||||
/* Test target enter data and target update from the target using map
|
||||
iterators. */
|
||||
|
||||
#include <stdlib.h>
|
||||
|
||||
#define DIM1 8
|
||||
#define DIM2 15
|
||||
|
||||
void mkarray (int *x[])
|
||||
{
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
{
|
||||
x[i] = (int *) malloc (DIM2 * sizeof (int));
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
x[i][j] = 0;
|
||||
}
|
||||
}
|
||||
|
||||
int main (void)
|
||||
{
|
||||
int *x[DIM1];
|
||||
int sum, expected;
|
||||
|
||||
mkarray (x);
|
||||
|
||||
#pragma omp target enter data map(alloc: x[:DIM1])
|
||||
#pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
|
||||
#pragma omp target map(from: expected)
|
||||
{
|
||||
expected = 0;
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
{
|
||||
x[i][j] = (i + 1) * (j + 2);
|
||||
expected += x[i][j];
|
||||
}
|
||||
}
|
||||
|
||||
/* Host copy of x should remain unchanged. */
|
||||
sum = 0;
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
sum += x[i][j];
|
||||
if (sum != 0)
|
||||
return 1;
|
||||
|
||||
#pragma omp target update from(iterator(i=0:DIM1): x[i][:DIM2])
|
||||
|
||||
/* Host copy should now be updated. */
|
||||
sum = 0;
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
sum += x[i][j];
|
||||
return sum - expected;
|
||||
}
|
||||
|
|
@ -0,0 +1,67 @@
|
|||
/* { dg-do run } */
|
||||
/* { dg-require-effective-target offload_device_nonshared_as } */
|
||||
|
||||
/* Test target enter data and target update to the target using map
|
||||
iterators with a function. */
|
||||
|
||||
#include <stdlib.h>
|
||||
|
||||
#define DIM1 8
|
||||
#define DIM2 15
|
||||
|
||||
void mkarray (int *x[])
|
||||
{
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
{
|
||||
x[i] = (int *) malloc (DIM2 * sizeof (int));
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
x[i][j] = rand ();
|
||||
}
|
||||
}
|
||||
|
||||
int f (int i)
|
||||
{
|
||||
return i * 2;
|
||||
}
|
||||
|
||||
int main (void)
|
||||
{
|
||||
int *x[DIM1], x_new[DIM1][DIM2];
|
||||
int sum, expected;
|
||||
|
||||
mkarray (x);
|
||||
|
||||
#pragma omp target enter data map(alloc: x[:DIM1])
|
||||
#pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
|
||||
|
||||
/* Update x on host. */
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
{
|
||||
x_new[i][j] = x[i][j];
|
||||
x[i][j] = (i + 1) * (j + 2);
|
||||
}
|
||||
|
||||
/* Update a subset of x on target. */
|
||||
#pragma omp target update to(iterator(i=0:DIM1/2): x[f (i)][:DIM2])
|
||||
|
||||
#pragma omp target map(from: sum)
|
||||
{
|
||||
sum = 0;
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
sum += x[i][j];
|
||||
}
|
||||
|
||||
/* Calculate expected value on host. */
|
||||
for (int i = 0; i < DIM1/2; i++)
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
x_new[f (i)][j] = x[f (i)][j];
|
||||
|
||||
expected = 0;
|
||||
for (int i = 0; i < DIM1; i++)
|
||||
for (int j = 0; j < DIM2; j++)
|
||||
expected += x_new[i][j];
|
||||
|
||||
return sum - expected;
|
||||
}
|
||||
Loading…
Reference in New Issue