2015-10-27 Cesar Philippidis <cesar@codesourcery.com>
gcc/c/
* c-parser.c (c_parser_oacc_shape_clause): Backport from trunk.
(c_parser_omp_simple_clause): Likewise.
(c_parser_oacc_all_clauses): Likewise.
gcc/cp/
* parser.c (cp_parser_oacc_shape_clause): Backport from trunk.
(cp_parser_oacc_all_clauses): Likewise.
* semantics.c (finish_omp_clauses): Likewise.
gcc/fortran/
* gfortran.h (gfc_omp_namelist): Add locus where member.
* openmp.c (gfc_free_omp_clauses): Recursively deallocate device_type
clauses.
(gfc_match_omp_variable_list): New function.
(resolve_omp_clauses): Remove where argument and use the where
gfc_omp_namespace member when reporting errors. Use
resolve_omp_duplicate_list to check for variables appearing in
mulitple clauses.
(gfc_match_omp_clauses): Update call to resolve_omp_clauses.
(gfc_match_oacc_declare): Likewise.
(resolve_omp_do): Likewise.
(resolve_oacc_loop): Likewise.
(gfc_resolve_oacc_directive): Likewise.
(gfc_resolve_omp_directive): Likewise.
(gfc_resolve_omp_declare_simd): Likewise.
(resolve_oacc_declare_map): New function.
(gfc_resolve_oacc_declare): Use it.
* trans-openmp.c (gfc_filter_oacc_combined_clauses): New function.
(gfc_trans_oacc_combined_directive): Use it.
gcc/testsuite/
* c-c++-common/goacc/loop-shape.c (int main): New test.
* g++.dg/gomp/pr33372-1.C: Adjust expected error messages.
* g++.dg/gomp/pr33372-3.C: Likewise.
* gfortran.dg/goacc/combined-directives.f90: New test.
* gfortran.dg/goacc/declare-2.f95: Adjust error message.
* gfortran.dg/goacc/multi-clause.f90: New test.
* gfortran.dg/gomp/intentin1.f90: Adjust error message.
libgomp/
* testsuite/libgomp.oacc-fortran/combdir-1.f90: Rename to ...
* testsuite/libgomp.oacc-fortran/combined-directive-1.f90: ... this.
Add a description of the test at the top of the file.
@@ -11226,119 +11226,146 @@ c_parser_omp_clause_is_device_ptr (c_parser *parser, tree list)
}
/* OpenACC:
- gang [( gang_expr_list )]
- worker [( expression )]
- vector [( expression )] */
+
+ gang [( gang-arg-list )]
+ worker [( [num:] int-expr )]
+ vector [( [length:] int-expr )]
+
+ where gang-arg is one of:
+
+ [num:] int-expr
+ static: size-expr
+
+ and size-expr may be:
+
+ *
+ int-expr
+*/
static tree
-c_parser_oacc_shape_clause (c_parser *parser, pragma_omp_clause c_kind,
+c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind,
const char *str, tree list)
{
- omp_clause_code kind;
const char *id = "num";
-
- switch (c_kind)
- {
- default:
- gcc_unreachable ();
- case PRAGMA_OACC_CLAUSE_GANG:
- kind = OMP_CLAUSE_GANG;
- break;
- case PRAGMA_OACC_CLAUSE_VECTOR:
- kind = OMP_CLAUSE_VECTOR;
- id = "length";
- break;
- case PRAGMA_OACC_CLAUSE_WORKER:
- kind = OMP_CLAUSE_WORKER;
- break;
- }
-
- tree op0 = NULL_TREE, op1 = NULL_TREE;
+ tree ops[2] = { NULL_TREE, NULL_TREE }, c;
location_t loc = c_parser_peek_token (parser)->location;
+ if (kind == OMP_CLAUSE_VECTOR)
+ id = "length";
+
if (c_parser_next_token_is (parser, CPP_OPEN_PAREN))
{
- tree *op_to_parse = &op0;
c_parser_consume_token (parser);
do
{
- if (c_parser_next_token_is (parser, CPP_NAME)
- || c_parser_next_token_is (parser, CPP_KEYWORD))
+ c_token *next = c_parser_peek_token (parser);
+ int idx = 0;
+
+ /* Gang static argument. */
+ if (kind == OMP_CLAUSE_GANG
+ && c_parser_next_token_is_keyword (parser, RID_STATIC))
{
- tree name_kind = c_parser_peek_token (parser)->value;
- const char *p = IDENTIFIER_POINTER (name_kind);
- if (kind == OMP_CLAUSE_GANG && strcmp ("static", p) == 0)
+ c_parser_consume_token (parser);
+
+ if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
+ goto cleanup_error;
+
+ idx = 1;
+ if (ops[idx] != NULL_TREE)
{
- c_parser_consume_token (parser);
- if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
- {
- c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
- return list;
- }
- op_to_parse = &op1;
- if (c_parser_next_token_is (parser, CPP_MULT))
- {
- c_parser_consume_token (parser);
- *op_to_parse = integer_minus_one_node;
- continue;
- }
+ c_parser_error (parser, "too many %<static%> arguments");
+ goto cleanup_error;
}
- else if (strcmp (id, p) == 0)
+
+ /* Check for the '*' argument. */
+ if (c_parser_next_token_is (parser, CPP_MULT))
{
c_parser_consume_token (parser);
- if (!c_parser_require (parser, CPP_COLON, "expected %<:%>"))
+ ops[idx] = integer_minus_one_node;
+
+ if (c_parser_next_token_is (parser, CPP_COMMA))
{
- c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
- return list;
+ c_parser_consume_token (parser);
+ continue;
}
- }
- else
- {
- if (kind == OMP_CLAUSE_GANG)
- c_parser_error (parser, "expected %<%num%> or %<static%>");
- else if (kind == OMP_CLAUSE_VECTOR)
- c_parser_error (parser, "expected %<length%>");
else
- c_parser_error (parser, "expected %<num%>");
- c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
- return list;
+ break;
}
}
+ /* Worker num: argument and vector length: arguments. */
+ else if (c_parser_next_token_is (parser, CPP_NAME)
+ && strcmp (id, IDENTIFIER_POINTER (next->value)) == 0
+ && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+ {
+ c_parser_consume_token (parser); /* id */
+ c_parser_consume_token (parser); /* ':' */
+ }
- if (*op_to_parse != NULL_TREE)
+ /* Now collect the actual argument. */
+ if (ops[idx] != NULL_TREE)
{
- c_parser_error (parser, "duplicate operand to clause");
- c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
- return list;
+ c_parser_error (parser, "unexpected argument");
+ goto cleanup_error;
}
location_t expr_loc = c_parser_peek_token (parser)->location;
- tree expr = c_parser_expression (parser).value;
+ tree expr = c_parser_expr_no_commas (parser, NULL).value;
if (expr == error_mark_node)
+ goto cleanup_error;
+
+ mark_exp_read (expr);
+ expr = c_fully_fold (expr, false, NULL);
+
+ /* Attempt to statically determine when the number isn't a
+ positive integer. */
+
+ if (!INTEGRAL_TYPE_P (TREE_TYPE (expr)))
{
- c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
+ c_parser_error (parser, "expected integer expression");
return list;
}
- mark_exp_read (expr);
- require_positive_expr (expr, expr_loc, str);
- *op_to_parse = expr;
- op_to_parse = &op0;
+ tree c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, expr,
+ build_int_cst (TREE_TYPE (expr), 0));
+ if (c == boolean_true_node)
+ {
+ warning_at (loc, 0,
+ "%<%s%> value must be positive", str);
+ expr = integer_one_node;
+ }
+
+ ops[idx] = expr;
+
+ if (kind == OMP_CLAUSE_GANG
+ && c_parser_next_token_is (parser, CPP_COMMA))
+ {
+ c_parser_consume_token (parser);
+ continue;
+ }
+ break;
}
- while (!c_parser_next_token_is (parser, CPP_CLOSE_PAREN));
- c_parser_consume_token (parser);
+ while (1);
+
+ if (!c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>"))
+ goto cleanup_error;
}
check_no_duplicate_clause (list, kind, str);
- tree c = build_omp_clause (loc, kind);
- if (op0)
- OMP_CLAUSE_OPERAND (c, 0) = op0;
- if (op1)
- OMP_CLAUSE_OPERAND (c, 1) = op1;
+ c = build_omp_clause (loc, kind);
+
+ if (ops[1])
+ OMP_CLAUSE_OPERAND (c, 1) = ops[1];
+
+ OMP_CLAUSE_OPERAND (c, 0) = ops[0];
OMP_CLAUSE_CHAIN (c) = list;
+
return c;
+
+ cleanup_error:
+ c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
+ return list; return c;
}
/* OpenACC:
@@ -11889,7 +11916,7 @@ c_parser_omp_clause_shared (c_parser *parser, tree list)
seq */
static tree
-c_parser_omp_simple_clause (c_parser *parser ATTRIBUTE_UNUSED,
+c_parser_omp_simple_clause (c_parser *parser,
enum omp_clause_code code, tree list)
{
check_no_duplicate_clause (list, code, omp_clause_code_name[code]);
@@ -12757,8 +12784,8 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
break;
case PRAGMA_OACC_CLAUSE_GANG:
c_name = "gang";
- clauses = c_parser_oacc_shape_clause (parser, c_kind, c_name,
- clauses);
+ clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG,
+ c_name, clauses);
break;
case PRAGMA_OACC_CLAUSE_HOST:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
@@ -12835,8 +12862,8 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
break;
case PRAGMA_OACC_CLAUSE_VECTOR:
c_name = "vector";
- clauses = c_parser_oacc_shape_clause (parser, c_kind, c_name,
- clauses);
+ clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
+ c_name, clauses);
break;
case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
c_name = "vector_length";
@@ -12849,8 +12876,8 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
break;
case PRAGMA_OACC_CLAUSE_WORKER:
c_name = "worker";
- clauses = c_parser_oacc_shape_clause (parser, c_kind, c_name,
- clauses);
+ clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_WORKER,
+ c_name, clauses);
break;
default:
c_parser_error (parser, "expected %<#pragma acc%> clause");
@@ -29712,142 +29712,126 @@ cp_parser_omp_positive_int_clause (cp_parser *parser, pragma_omp_clause c_kind,
}
/* OpenACC:
- gang [( gang_expr_list )]
- worker [( expression )]
- vector [( expression )] */
+
+ gang [( gang-arg-list )]
+ worker [( [num:] int-expr )]
+ vector [( [length:] int-expr )]
+
+ where gang-arg is one of:
+
+ [num:] int-expr
+ static: size-expr
+
+ and size-expr may be:
+
+ *
+ int-expr
+*/
static tree
-cp_parser_oacc_shape_clause (cp_parser *parser, pragma_omp_clause c_kind,
+cp_parser_oacc_shape_clause (cp_parser *parser, omp_clause_code kind,
const char *str, tree list)
{
- omp_clause_code kind;
const char *id = "num";
cp_lexer *lexer = parser->lexer;
-
- switch (c_kind)
- {
- default:
- gcc_unreachable ();
- case PRAGMA_OACC_CLAUSE_GANG:
- kind = OMP_CLAUSE_GANG;
- break;
- case PRAGMA_OACC_CLAUSE_VECTOR:
- kind = OMP_CLAUSE_VECTOR;
- id = "length";
- break;
- case PRAGMA_OACC_CLAUSE_WORKER:
- kind = OMP_CLAUSE_WORKER;
- break;
- }
-
- tree op0 = NULL_TREE, op1 = NULL_TREE;
+ tree ops[2] = { NULL_TREE, NULL_TREE }, c;
location_t loc = cp_lexer_peek_token (lexer)->location;
+ if (kind == OMP_CLAUSE_VECTOR)
+ id = "length";
+
if (cp_lexer_next_token_is (lexer, CPP_OPEN_PAREN))
{
- tree *op_to_parse = &op0;
cp_lexer_consume_token (lexer);
do
{
- if (cp_lexer_next_token_is (lexer, CPP_NAME)
- || cp_lexer_next_token_is (lexer, CPP_KEYWORD))
+ cp_token *next = cp_lexer_peek_token (lexer);
+ int idx = 0;
+
+ /* Gang static argument. */
+ if (kind == OMP_CLAUSE_GANG
+ && cp_lexer_next_token_is_keyword (lexer, RID_STATIC))
{
- tree name_kind = cp_lexer_peek_token (lexer)->u.value;
- const char *p = IDENTIFIER_POINTER (name_kind);
- if (kind == OMP_CLAUSE_GANG && strcmp ("static", p) == 0)
+ cp_lexer_consume_token (lexer);
+
+ if (!cp_parser_require (parser, CPP_COLON, RT_COLON))
+ goto cleanup_error;
+
+ idx = 1;
+ if (ops[idx] != NULL)
{
- cp_lexer_consume_token (lexer);
- if (!cp_parser_require (parser, CPP_COLON, RT_COLON))
- {
- cp_parser_skip_to_closing_parenthesis (parser, false,
- false, true);
- return list;
- }
- op_to_parse = &op1;
- if (cp_lexer_next_token_is (lexer, CPP_MULT))
- {
- if (*op_to_parse != NULL_TREE)
- {
- cp_parser_error (parser,
- "duplicate %<num%> argument");
- cp_parser_skip_to_closing_parenthesis (parser,
- false, false,
- true);
- return list;
- }
- cp_lexer_consume_token (lexer);
- *op_to_parse = integer_minus_one_node;
- if (cp_lexer_next_token_is (lexer, CPP_COMMA))
- cp_lexer_consume_token (lexer);
- continue;
- }
+ cp_parser_error (parser, "too many %<static%> arguments");
+ goto cleanup_error;
}
- else if (strcmp (id, p) == 0)
+
+ /* Check for the '*' argument. */
+ if (cp_lexer_next_token_is (lexer, CPP_MULT))
{
- op_to_parse = &op0;
cp_lexer_consume_token (lexer);
- if (!cp_parser_require (parser, CPP_COLON, RT_COLON))
+ ops[idx] = integer_minus_one_node;
+
+ if (cp_lexer_next_token_is (lexer, CPP_COMMA))
{
- cp_parser_skip_to_closing_parenthesis (parser, false,
- false, true);
- return list;
+ cp_lexer_consume_token (lexer);
+ continue;
}
- }
- else
- {
- if (kind == OMP_CLAUSE_GANG)
- cp_parser_error (parser,
- "expected %<%num%> or %<static%>");
- else if (kind == OMP_CLAUSE_VECTOR)
- cp_parser_error (parser, "expected %<length%>");
- else
- cp_parser_error (parser, "expected %<num%>");
- cp_parser_skip_to_closing_parenthesis (parser, false, false,
- true);
- return list;
+ else break;
}
}
+ /* Worker num: argument and vector length: arguments. */
+ else if (cp_lexer_next_token_is (lexer, CPP_NAME)
+ && strcmp (id, IDENTIFIER_POINTER (next->u.value)) == 0
+ && cp_lexer_nth_token_is (lexer, 2, CPP_COLON))
+ {
+ cp_lexer_consume_token (lexer); /* id */
+ cp_lexer_consume_token (lexer); /* ':' */
+ }
- if (*op_to_parse != NULL_TREE)
+ /* Now collect the actual argument. */
+ if (ops[idx] != NULL_TREE)
{
- cp_parser_error (parser, "duplicate operand to clause");
- cp_parser_skip_to_closing_parenthesis (parser, false, false,
- true);
- return list;
+ cp_parser_error (parser, "unexpected argument");
+ goto cleanup_error;
}
- location_t expr_loc = cp_lexer_peek_token (lexer)->location;
tree expr = cp_parser_assignment_expression (parser, NULL, false,
false);
if (expr == error_mark_node)
- {
- cp_parser_skip_to_closing_parenthesis (parser, false, false,
- true);
- return list;
- }
+ goto cleanup_error;
mark_exp_read (expr);
- require_positive_expr (expr, expr_loc, str);
- *op_to_parse = expr;
- op_to_parse = &op0;
+ ops[idx] = expr;
- if (cp_lexer_next_token_is (lexer, CPP_COMMA))
- cp_lexer_consume_token (lexer);
+ if (kind == OMP_CLAUSE_GANG
+ && cp_lexer_next_token_is (lexer, CPP_COMMA))
+ {
+ cp_lexer_consume_token (lexer);
+ continue;
+ }
+ break;
}
- while (!cp_lexer_next_token_is (lexer, CPP_CLOSE_PAREN));
- cp_lexer_consume_token (lexer);
+ while (1);
+
+ if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN))
+ goto cleanup_error;
}
check_no_duplicate_clause (list, kind, str, loc);
- tree c = build_omp_clause (loc, kind);
- if (op0)
- OMP_CLAUSE_OPERAND (c, 0) = op0;
- if (op1)
- OMP_CLAUSE_OPERAND (c, 1) = op1;
+ c = build_omp_clause (loc, kind);
+
+ if (ops[1])
+ OMP_CLAUSE_OPERAND (c, 1) = ops[1];
+
+ OMP_CLAUSE_OPERAND (c, 0) = ops[0];
OMP_CLAUSE_CHAIN (c) = list;
+
return c;
+
+ cleanup_error:
+ cp_parser_skip_to_closing_parenthesis (parser, false, false, true);
+ return list;
}
/* OpenACC 2.0:
@@ -31712,8 +31696,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
break;
case PRAGMA_OACC_CLAUSE_GANG:
c_name = "gang";
- clauses = cp_parser_oacc_shape_clause (parser, c_kind, c_name,
- clauses);
+ clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG,
+ c_name, clauses);
break;
case PRAGMA_OACC_CLAUSE_HOST:
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
@@ -31783,8 +31767,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
break;
case PRAGMA_OACC_CLAUSE_VECTOR:
c_name = "vector";
- clauses = cp_parser_oacc_shape_clause (parser, c_kind, c_name,
- clauses);
+ clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
+ c_name, clauses);
break;
case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
c_name = "vector_length";
@@ -31797,8 +31781,8 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
break;
case PRAGMA_OACC_CLAUSE_WORKER:
c_name = "worker";
- clauses = cp_parser_oacc_shape_clause (parser, c_kind, c_name,
- clauses);
+ clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_WORKER,
+ c_name, clauses);
break;
default:
cp_parser_error (parser, "expected %<#pragma acc%> clause");
@@ -5986,37 +5986,6 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
bitmap_set_bit (&firstprivate_head, DECL_UID (t));
goto handle_field_decl;
- case OMP_CLAUSE_GANG:
- case OMP_CLAUSE_VECTOR:
- case OMP_CLAUSE_WORKER:
- /* Operand 0 is the num: or length: argument. */
- t = OMP_CLAUSE_OPERAND (c, 0);
- if (t == NULL_TREE)
- break;
-
- t = maybe_convert_cond (t);
- if (t == error_mark_node)
- remove = true;
- else if (!processing_template_decl)
- t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
- OMP_CLAUSE_OPERAND (c, 0) = t;
-
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_GANG)
- break;
-
- /* Ooperand 1 is the gang static: argument. */
- t = OMP_CLAUSE_OPERAND (c, 1);
- if (t == NULL_TREE)
- break;
-
- t = maybe_convert_cond (t);
- if (t == error_mark_node)
- remove = true;
- else if (!processing_template_decl)
- t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
- OMP_CLAUSE_OPERAND (c, 1) = t;
- break;
-
case OMP_CLAUSE_LASTPRIVATE:
t = omp_clause_decl_field (OMP_CLAUSE_DECL (c));
if (t)
@@ -6071,6 +6040,48 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
OMP_CLAUSE_FINAL_EXPR (c) = t;
break;
+ case OMP_CLAUSE_GANG:
+ /* Operand 1 is the gang static: argument. */
+ t = OMP_CLAUSE_OPERAND (c, 1);
+ if (t != NULL_TREE)
+ {
+ if (t == error_mark_node)
+ remove = true;
+ else if (!type_dependent_expression_p (t)
+ && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
+ {
+ error ("%<gang%> static expression must be integral");
+ remove = true;
+ }
+ else
+ {
+ t = mark_rvalue_use (t);
+ if (!processing_template_decl)
+ {
+ t = maybe_constant_value (t);
+ if (TREE_CODE (t) == INTEGER_CST
+ && tree_int_cst_sgn (t) != 1
+ && t != integer_minus_one_node)
+ {
+ warning_at (OMP_CLAUSE_LOCATION (c), 0,
+ "%<gang%> static value must be"
+ "positive");
+ t = integer_one_node;
+ }
+ }
+ t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
+ }
+ OMP_CLAUSE_OPERAND (c, 1) = t;
+ }
+ /* Check operand 0, the num argument. */
+
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ if (OMP_CLAUSE_OPERAND (c, 0) == NULL_TREE)
+ break;
+
+ case OMP_CLAUSE_NUM_TASKS:
+ case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_GANGS:
case OMP_CLAUSE_NUM_WORKERS:
@@ -6083,18 +6094,21 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
{
switch (OMP_CLAUSE_CODE (c))
{
- case OMP_CLAUSE_NUM_THREADS:
- error ("num_threads expression must be integral"); break;
- case OMP_CLAUSE_NUM_GANGS:
- error ("%<num_gangs%> expression must be integral"); break;
- case OMP_CLAUSE_NUM_WORKERS:
- error ("%<num_workers%> expression must be integral");
+ case OMP_CLAUSE_GANG:
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<gang%> num expression must be integral"); break;
+ case OMP_CLAUSE_VECTOR:
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<vector%> length expression must be integral");
break;
- case OMP_CLAUSE_VECTOR_LENGTH:
- error ("%<vector_length%> expression must be integral");
+ case OMP_CLAUSE_WORKER:
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<worker%> num expression must be integral");
break;
default:
- error ("invalid argument");
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qs expression must be integral",
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
}
remove = true;
}
@@ -6107,9 +6121,28 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
if (TREE_CODE (t) == INTEGER_CST
&& tree_int_cst_sgn (t) != 1)
{
- warning_at (OMP_CLAUSE_LOCATION (c), 0,
- /* TODO */
- "%<num_threads%> value must be positive");
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_GANG:
+ warning_at (OMP_CLAUSE_LOCATION (c), 0,
+ "%<gang%> num value must be positive");
+ break;
+ case OMP_CLAUSE_VECTOR:
+ warning_at (OMP_CLAUSE_LOCATION (c), 0,
+ "%<vector%> length value must be"
+ "positive");
+ break;
+ case OMP_CLAUSE_WORKER:
+ warning_at (OMP_CLAUSE_LOCATION (c), 0,
+ "%<worker%> num value must be"
+ "positive");
+ break;
+ default:
+ warning_at (OMP_CLAUSE_LOCATION (c), 0,
+ "%qs value must be positive",
+ omp_clause_code_name
+ [OMP_CLAUSE_CODE (c)]);
+ }
t = integer_one_node;
}
t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
@@ -6186,35 +6219,6 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
}
break;
- case OMP_CLAUSE_NUM_TEAMS:
- t = OMP_CLAUSE_NUM_TEAMS_EXPR (c);
- if (t == error_mark_node)
- remove = true;
- else if (!type_dependent_expression_p (t)
- && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
- {
- error ("%<num_teams%> expression must be integral");
- remove = true;
- }
- else
- {
- t = mark_rvalue_use (t);
- if (!processing_template_decl)
- {
- t = maybe_constant_value (t);
- if (TREE_CODE (t) == INTEGER_CST
- && tree_int_cst_sgn (t) != 1)
- {
- warning_at (OMP_CLAUSE_LOCATION (c), 0,
- "%<num_teams%> value must be positive");
- t = integer_one_node;
- }
- t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
- }
- OMP_CLAUSE_NUM_TEAMS_EXPR (c) = t;
- }
- break;
-
case OMP_CLAUSE_ASYNC:
t = OMP_CLAUSE_ASYNC_EXPR (c);
if (t == error_mark_node)
@@ -6667,35 +6671,6 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
}
goto check_dup_generic;
- case OMP_CLAUSE_NUM_TASKS:
- t = OMP_CLAUSE_NUM_TASKS_EXPR (c);
- if (t == error_mark_node)
- remove = true;
- else if (!type_dependent_expression_p (t)
- && !INTEGRAL_TYPE_P (TREE_TYPE (t)))
- {
- error ("%<num_tasks%> expression must be integral");
- remove = true;
- }
- else
- {
- t = mark_rvalue_use (t);
- if (!processing_template_decl)
- {
- t = maybe_constant_value (t);
- if (TREE_CODE (t) == INTEGER_CST
- && tree_int_cst_sgn (t) != 1)
- {
- warning_at (OMP_CLAUSE_LOCATION (c), 0,
- "%<num_tasks%> value must be positive");
- t = integer_one_node;
- }
- t = fold_build_cleanup_point_expr (TREE_TYPE (t), t);
- }
- OMP_CLAUSE_NUM_TASKS_EXPR (c) = t;
- }
- break;
-
case OMP_CLAUSE_GRAINSIZE:
t = OMP_CLAUSE_GRAINSIZE_EXPR (c);
if (t == error_mark_node)
@@ -1136,6 +1136,7 @@ typedef struct gfc_omp_namelist
} u;
struct gfc_omp_namelist_udr *udr;
struct gfc_omp_namelist *next;
+ locus where;
}
gfc_omp_namelist;
@@ -87,6 +87,7 @@ gfc_free_omp_clauses (gfc_omp_clauses *c)
gfc_free_omp_namelist (c->lists[i]);
gfc_free_expr_list (c->wait_list);
gfc_free_expr_list (c->tile_list);
+ gfc_free_omp_clauses (c->dtype_clauses);
free (c);
}
@@ -263,6 +264,7 @@ gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
}
tail->sym = sym;
tail->expr = expr;
+ tail->where = cur_loc;
goto next_item;
case MATCH_NO:
break;
@@ -297,6 +299,7 @@ gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
tail = tail->next;
}
tail->sym = sym;
+ tail->where = cur_loc;
}
next_item:
@@ -1249,14 +1252,6 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, uint64_t mask,
if (gfc_match_omp_eos () != MATCH_YES)
{
- gfc_omp_clauses *t;
- c = base_clauses->dtype_clauses;
- while (c)
- {
- t = c->dtype_clauses;
- gfc_free_omp_clauses (c);
- c = t;
- }
gfc_free_omp_clauses (base_clauses);
return MATCH_ERROR;
}
@@ -1473,8 +1468,8 @@ gfc_match_oacc_declare (void)
if (n->u.map_op != OMP_MAP_FORCE_ALLOC
&& n->u.map_op != OMP_MAP_FORCE_TO)
{
- gfc_error ("Invalid clause in module with "
- "$!ACC DECLARE at %C");
+ gfc_error ("Invalid clause in module with $!ACC DECLARE at %L",
+ &n->where);
return MATCH_ERROR;
}
@@ -1483,29 +1478,29 @@ gfc_match_oacc_declare (void)
if (ns->proc_name->attr.oacc_function)
{
- gfc_error ("Invalid declare in routine with " "$!ACC DECLARE at %C");
+ gfc_error ("Invalid declare in routine with $!ACC DECLARE at %C");
return MATCH_ERROR;
}
if (s->attr.in_common)
{
- gfc_error ("Unsupported: variable in a common block with "
- "$!ACC DECLARE at %C");
+ gfc_error ("Variable in a common block with $!ACC DECLARE at %L",
+ &n->where);
return MATCH_ERROR;
}
if (s->attr.use_assoc)
{
- gfc_error ("Unsupported: variable is USE-associated with "
- "$!ACC DECLARE at %C");
+ gfc_error ("Variable is USE-associated with $!ACC DECLARE at %L",
+ &n->where);
return MATCH_ERROR;
}
if ((s->attr.dimension || s->attr.codimension)
&& s->attr.dummy && s->as->type != AS_EXPLICIT)
{
- gfc_error ("Unsupported: assumed-size dummy array with "
- "$!ACC DECLARE at %C");
+ gfc_error ("Assumed-size dummy array with $!ACC DECLARE at %L",
+ &n->where);
return MATCH_ERROR;
}
@@ -1533,37 +1528,6 @@ gfc_match_oacc_declare (void)
new_oc->module_var = module_var;
new_oc->clauses = c;
new_oc->where = gfc_current_locus;
-
- for (oc = new_oc; oc; oc = oc->next)
- {
- c = oc->clauses;
- for (n = c->lists[OMP_LIST_MAP]; n != NULL; n = n->next)
- n->sym->mark = 0;
- }
-
- for (oc = new_oc; oc; oc = oc->next)
- {
- c = oc->clauses;
- for (n = c->lists[OMP_LIST_MAP]; n != NULL; n = n->next)
- {
- if (n->sym->mark)
- {
- gfc_error ("Symbol %qs present on multiple clauses at %C",
- n->sym->name);
- return MATCH_ERROR;
- }
- else
- n->sym->mark = 1;
- }
- }
-
- for (oc = new_oc; oc; oc = oc->next)
- {
- c = oc->clauses;
- for (n = c->lists[OMP_LIST_MAP]; n != NULL; n = n->next)
- n->sym->mark = 1;
- }
-
ns->oacc_declare = new_oc;
return MATCH_YES;
@@ -3187,36 +3151,47 @@ resolve_omp_udr_clause (gfc_omp_namelist *n, gfc_namespace *ns,
return copy;
}
-/* Returns true if clause in list 'list' is compatible with any of
- of the clauses in lists [0..list-1]. E.g., a reduction variable may
- appear in both reduction and private clauses, so this function
- will return true in this case. */
+/* Check if a variable appears in multiple clauses. */
-static bool
-oacc_compatible_clauses (gfc_omp_clauses *clauses, int list,
- gfc_symbol *sym, bool openacc)
+static void
+resolve_omp_duplicate_list (gfc_omp_namelist *clause_list, bool openacc,
+ int list)
{
gfc_omp_namelist *n;
+ const char *error_msg = "Symbol %qs present on multiple clauses at %L";
- if (!openacc)
- return false;
+ /* OpenACC reduction clauses are compatible with everything. We only
+ need to check if a reduction variable is used more than once. */
+ if (openacc && list == OMP_LIST_REDUCTION)
+ {
+ hash_set<gfc_symbol *> reductions;
- if (list != OMP_LIST_REDUCTION)
- return false;
+ for (n = clause_list; n; n = n->next)
+ {
+ if (reductions.contains (n->sym))
+ gfc_error (error_msg, n->sym->name, &n->where);
+ else
+ reductions.add (n->sym);
+ }
- for (n = clauses->lists[OMP_LIST_FIRST]; n; n = n->next)
- if (n->sym == sym)
- return true;
+ return;
+ }
- return false;
+ /* Ensure that variables are only used in one clause. */
+ for (n = clause_list; n; n = n->next)
+ {
+ if (n->sym->mark)
+ gfc_error (error_msg, n->sym->name, &n->where);
+ else
+ n->sym->mark = 1;
+ }
}
/* OpenMP directive resolving routines. */
static void
-resolve_omp_clauses (gfc_code *code, locus *where,
- gfc_omp_clauses *omp_clauses, gfc_namespace *ns,
- bool openacc = false)
+resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
+ gfc_namespace *ns, bool openacc = false)
{
gfc_omp_namelist *n;
gfc_expr_list *el;
@@ -3275,7 +3250,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
{
if (!code && (!n->sym->attr.dummy || n->sym->ns != ns))
gfc_error ("Variable %qs is not a dummy argument at %L",
- n->sym->name, where);
+ n->sym->name, &n->where);
continue;
}
if (n->sym->attr.flavor == FL_PROCEDURE
@@ -3307,7 +3282,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
}
}
gfc_error ("Object %qs is not a variable at %L", n->sym->name,
- where);
+ &n->where);
}
for (list = 0; list < OMP_LIST_NUM; list++)
@@ -3318,57 +3293,22 @@ resolve_omp_clauses (gfc_code *code, locus *where,
&& (list != OMP_LIST_MAP || openacc)
&& list != OMP_LIST_FROM
&& list != OMP_LIST_TO)
- for (n = omp_clauses->lists[list]; n; n = n->next)
- {
- if (n->sym->mark && !oacc_compatible_clauses (omp_clauses, list,
- n->sym, openacc))
- gfc_error ("Symbol %qs present on multiple clauses at %L",
- n->sym->name, where);
- else
- n->sym->mark = 1;
- }
+ resolve_omp_duplicate_list (omp_clauses->lists[list], openacc, list);
- gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1);
- for (list = OMP_LIST_FIRSTPRIVATE; list <= OMP_LIST_LASTPRIVATE; list++)
- for (n = omp_clauses->lists[list]; n; n = n->next)
- if (n->sym->mark)
- {
- gfc_error ("Symbol %qs present on multiple clauses at %L",
- n->sym->name, where);
- n->sym->mark = 0;
- }
+ resolve_omp_duplicate_list (omp_clauses->lists[OMP_LIST_FIRSTPRIVATE],
+ false, OMP_LIST_FIRSTPRIVATE);
- for (n = omp_clauses->lists[OMP_LIST_FIRSTPRIVATE]; n; n = n->next)
- {
- if (n->sym->mark)
- gfc_error ("Symbol %qs present on multiple clauses at %L",
- n->sym->name, where);
- else
- n->sym->mark = 1;
- }
for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next)
n->sym->mark = 0;
- for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next)
- {
- if (n->sym->mark)
- gfc_error ("Symbol %qs present on multiple clauses at %L",
- n->sym->name, where);
- else
- n->sym->mark = 1;
- }
+ resolve_omp_duplicate_list (omp_clauses->lists[OMP_LIST_LASTPRIVATE],
+ false, OMP_LIST_LASTPRIVATE);
for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next)
n->sym->mark = 0;
- for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next)
- {
- if (n->sym->mark)
- gfc_error ("Symbol %qs present on multiple clauses at %L",
- n->sym->name, where);
- else
- n->sym->mark = 1;
- }
+ resolve_omp_duplicate_list (omp_clauses->lists[OMP_LIST_ALIGNED],
+ false, OMP_LIST_ALIGNED);
for (n = omp_clauses->lists[OMP_LIST_TO]; n; n = n->next)
n->sym->mark = 0;
@@ -3379,7 +3319,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
{
if (n->expr == NULL && n->sym->mark)
gfc_error ("Symbol %qs present on both FROM and TO clauses at %L",
- n->sym->name, where);
+ n->sym->name, &n->where);
else
n->sym->mark = 1;
}
@@ -3401,7 +3341,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
{
if (!n->sym->attr.threadprivate)
gfc_error ("Non-THREADPRIVATE object %qs in COPYIN clause"
- " at %L", n->sym->name, where);
+ " at %L", n->sym->name, &n->where);
}
break;
case OMP_LIST_COPYPRIVATE:
@@ -3409,10 +3349,10 @@ resolve_omp_clauses (gfc_code *code, locus *where,
{
if (n->sym->as && n->sym->as->type == AS_ASSUMED_SIZE)
gfc_error ("Assumed size array %qs in COPYPRIVATE clause "
- "at %L", n->sym->name, where);
+ "at %L", n->sym->name, &n->where);
if (n->sym->attr.pointer && n->sym->attr.intent == INTENT_IN)
gfc_error ("INTENT(IN) POINTER %qs in COPYPRIVATE clause "
- "at %L", n->sym->name, where);
+ "at %L", n->sym->name, &n->where);
}
break;
case OMP_LIST_SHARED:
@@ -3420,13 +3360,13 @@ resolve_omp_clauses (gfc_code *code, locus *where,
{
if (n->sym->attr.threadprivate)
gfc_error ("THREADPRIVATE object %qs in SHARED clause at "
- "%L", n->sym->name, where);
+ "%L", n->sym->name, &n->where);
if (n->sym->attr.cray_pointee)
gfc_error ("Cray pointee %qs in SHARED clause at %L",
- n->sym->name, where);
+ n->sym->name, &n->where);
if (n->sym->attr.associate_var)
gfc_error ("ASSOCIATE name %qs in SHARED clause at %L",
- n->sym->name, where);
+ n->sym->name, &n->where);
}
break;
case OMP_LIST_ALIGNED:
@@ -3442,7 +3382,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
!= ISOCBINDING_PTR)))
gfc_error ("%qs in ALIGNED clause must be POINTER, "
"ALLOCATABLE, Cray pointer or C_PTR at %L",
- n->sym->name, where);
+ n->sym->name, &n->where);
else if (n->expr)
{
gfc_expr *expr = n->expr;
@@ -3454,7 +3394,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
|| alignment <= 0)
gfc_error ("%qs in ALIGNED clause at %L requires a scalar "
"positive constant integer alignment "
- "expression", n->sym->name, where);
+ "expression", n->sym->name, &n->where);
}
}
break;
@@ -3472,10 +3412,11 @@ resolve_omp_clauses (gfc_code *code, locus *where,
|| n->expr->ref->next
|| n->expr->ref->type != REF_ARRAY)
gfc_error ("%qs in %s clause at %L is not a proper "
- "array section", n->sym->name, name, where);
+ "array section", n->sym->name, name,
+ &n->where);
else if (n->expr->ref->u.ar.codimen)
gfc_error ("Coarrays not supported in %s clause at %L",
- name, where);
+ name, &n->where);
else
{
int i;
@@ -3485,7 +3426,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
{
gfc_error ("Stride should not be specified for "
"array section in %s clause at %L",
- name, where);
+ name, &n->where);
break;
}
else if (ar->dimen_type[i] != DIMEN_ELEMENT
@@ -3493,7 +3434,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
{
gfc_error ("%qs in %s clause at %L is not a "
"proper array section",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
break;
}
else if (list == OMP_LIST_DEPEND
@@ -3506,7 +3447,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
{
gfc_error ("%qs in DEPEND clause at %L is a "
"zero size array section",
- n->sym->name, where);
+ n->sym->name, &n->where);
break;
}
}
@@ -3515,9 +3456,9 @@ resolve_omp_clauses (gfc_code *code, locus *where,
{
if (list == OMP_LIST_MAP
&& n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
- resolve_oacc_deviceptr_clause (n->sym, *where, name);
+ resolve_oacc_deviceptr_clause (n->sym, n->where, name);
else
- resolve_oacc_data_clauses (n->sym, *where, name);
+ resolve_oacc_data_clauses (n->sym, n->where, name);
}
}
@@ -3527,10 +3468,10 @@ resolve_omp_clauses (gfc_code *code, locus *where,
n->sym->attr.referenced = 1;
if (n->sym->attr.threadprivate)
gfc_error ("THREADPRIVATE object %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (n->sym->attr.cray_pointee)
gfc_error ("Cray pointee %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
}
break;
default:
@@ -3539,35 +3480,35 @@ resolve_omp_clauses (gfc_code *code, locus *where,
bool bad = false;
if (n->sym->attr.threadprivate)
gfc_error ("THREADPRIVATE object %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (n->sym->attr.cray_pointee)
gfc_error ("Cray pointee %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (n->sym->attr.associate_var)
gfc_error ("ASSOCIATE name %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (list != OMP_LIST_PRIVATE)
{
if (n->sym->attr.proc_pointer && list == OMP_LIST_REDUCTION)
gfc_error ("Procedure pointer %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (n->sym->attr.pointer && list == OMP_LIST_REDUCTION)
gfc_error ("POINTER object %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (n->sym->attr.cray_pointer && list == OMP_LIST_REDUCTION)
gfc_error ("Cray pointer %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
}
if (code
&& (oacc_is_loop (code) || code->op == EXEC_OACC_PARALLEL))
- check_array_not_assumed (n->sym, *where, name);
+ check_array_not_assumed (n->sym, n->where, name);
else if (n->sym->as && n->sym->as->type == AS_ASSUMED_SIZE)
gfc_error ("Assumed size array %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (n->sym->attr.in_namelist && list != OMP_LIST_REDUCTION)
gfc_error ("Variable %qs in %s clause is used in "
"NAMELIST statement at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (n->sym->attr.pointer && n->sym->attr.intent == INTENT_IN)
switch (list)
{
@@ -3576,7 +3517,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
case OMP_LIST_LINEAR:
/* case OMP_LIST_REDUCTION: */
gfc_error ("INTENT(IN) POINTER %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
break;
default:
break;
@@ -3670,7 +3611,7 @@ resolve_omp_clauses (gfc_code *code, locus *where,
}
gfc_error ("!$OMP DECLARE REDUCTION %s not found "
"for type %s at %L", udr_name,
- gfc_typename (&n->sym->ts), where);
+ gfc_typename (&n->sym->ts), &n->where);
}
else
{
@@ -3692,10 +3633,10 @@ resolve_omp_clauses (gfc_code *code, locus *where,
case OMP_LIST_LINEAR:
if (n->sym->ts.type != BT_INTEGER)
gfc_error ("LINEAR variable %qs must be INTEGER "
- "at %L", n->sym->name, where);
+ "at %L", n->sym->name, &n->where);
else if (!code && !n->sym->attr.value)
gfc_error ("LINEAR dummy argument %qs must have VALUE "
- "attribute at %L", n->sym->name, where);
+ "attribute at %L", n->sym->name, &n->where);
else if (n->expr)
{
gfc_expr *expr = n->expr;
@@ -3704,11 +3645,11 @@ resolve_omp_clauses (gfc_code *code, locus *where,
|| expr->rank != 0)
gfc_error ("%qs in LINEAR clause at %L requires "
"a scalar integer linear-step expression",
- n->sym->name, where);
+ n->sym->name, &n->where);
else if (!code && expr->expr_type != EXPR_CONSTANT)
gfc_error ("%qs in LINEAR clause at %L requires "
"a constant integer linear-step expression",
- n->sym->name, where);
+ n->sym->name, &n->where);
}
break;
/* Workaround for PR middle-end/26316, nothing really needs
@@ -3721,23 +3662,23 @@ resolve_omp_clauses (gfc_code *code, locus *where,
|| (n->sym->ts.type == BT_CLASS && CLASS_DATA (n->sym)
&& CLASS_DATA (n->sym)->attr.allocatable))
gfc_error ("ALLOCATABLE object %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (n->sym->attr.pointer
|| (n->sym->ts.type == BT_CLASS && CLASS_DATA (n->sym)
&& CLASS_DATA (n->sym)->attr.class_pointer))
gfc_error ("POINTER object %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (n->sym->attr.cray_pointer)
gfc_error ("Cray pointer object %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
if (n->sym->attr.cray_pointee)
gfc_error ("Cray pointee object %qs in %s clause at %L",
- n->sym->name, name, where);
+ n->sym->name, name, &n->where);
/* FALLTHRU */
case OMP_LIST_DEVICE_RESIDENT:
case OMP_LIST_CACHE:
- check_symbol_not_pointer (n->sym, *where, name);
- check_array_not_assumed (n->sym, *where, name);
+ check_symbol_not_pointer (n->sym, n->where, name);
+ check_array_not_assumed (n->sym, n->where, name);
break;
default:
break;
@@ -4503,7 +4444,7 @@ resolve_omp_do (gfc_code *code)
}
if (code->ext.omp_clauses)
- resolve_omp_clauses (code, &code->loc, code->ext.omp_clauses, NULL);
+ resolve_omp_clauses (code, code->ext.omp_clauses, NULL);
do_code = code->block->next;
collapse = code->ext.omp_clauses->collapse;
@@ -4940,7 +4881,7 @@ resolve_oacc_loop (gfc_code *code)
int collapse;
if (code->ext.omp_clauses)
- resolve_omp_clauses (code, &code->loc, code->ext.omp_clauses, NULL, true);
+ resolve_omp_clauses (code, code->ext.omp_clauses, NULL, true);
do_code = code->block->next;
collapse = code->ext.omp_clauses->collapse;
@@ -4950,6 +4891,26 @@ resolve_oacc_loop (gfc_code *code)
resolve_oacc_nested_loops (code, do_code, collapse, "collapsed");
}
+/* Helper function for gfc_resolve_oacc_declare. Scan omp_map_list LIST
+ in DECLARE at location LOC. */
+
+static void
+resolve_oacc_declare_map (gfc_oacc_declare *declare, int list)
+{
+ gfc_oacc_declare *oc;
+ gfc_omp_namelist *n;
+
+ for (oc = declare; oc; oc = oc->next)
+ for (n = oc->clauses->lists[list]; n; n = n->next)
+ n->sym->mark = 0;
+
+ for (oc = declare; oc; oc = oc->next)
+ resolve_omp_duplicate_list (oc->clauses->lists[list], false, list);
+
+ for (oc = declare; oc; oc = oc->next)
+ for (n = oc->clauses->lists[list]; n; n = n->next)
+ n->sym->mark = 0;
+}
void
gfc_resolve_oacc_declare (gfc_namespace *ns)
@@ -4966,64 +4927,27 @@ gfc_resolve_oacc_declare (gfc_namespace *ns)
{
loc = oc->where;
- for (list = OMP_LIST_DEVICE_RESIDENT;
- list <= OMP_LIST_DEVICE_RESIDENT; list++)
- for (n = oc->clauses->lists[list]; n; n = n->next)
- {
- n->sym->mark = 0;
- if (n->sym->attr.flavor == FL_PARAMETER)
- gfc_error ("PARAMETER object %qs is not allowed at %L",
- n->sym->name, &loc);
- }
-
- for (list = OMP_LIST_DEVICE_RESIDENT;
- list <= OMP_LIST_DEVICE_RESIDENT; list++)
- for (n = oc->clauses->lists[list]; n; n = n->next)
- {
- if (n->sym->mark)
- gfc_error ("Symbol %qs present on multiple clauses at %L",
- n->sym->name, &loc);
- else
- n->sym->mark = 1;
- }
-
for (n = oc->clauses->lists[OMP_LIST_DEVICE_RESIDENT]; n; n = n->next)
- check_array_not_assumed (n->sym, loc, "DEVICE_RESIDENT");
-
- for (n = oc->clauses->lists[OMP_LIST_MAP]; n; n = n->next)
{
- if (n->expr && n->expr->ref->type == REF_ARRAY)
- gfc_error ("Subarray: %qs not allowed in $!ACC DECLARE at %L",
- n->sym->name, &loc);
- }
- }
-
- for (oc = ns->oacc_declare; oc; oc = oc->next)
- {
- for (list = OMP_LIST_LINK; list <= OMP_LIST_LINK; list++)
- for (n = oc->clauses->lists[list]; n; n = n->next)
n->sym->mark = 0;
- }
+ if (n->sym->attr.flavor == FL_PARAMETER)
+ gfc_error ("PARAMETER object %qs is not allowed at %L",
+ n->sym->name, &n->where);
- for (oc = ns->oacc_declare; oc; oc = oc->next)
- {
- for (list = OMP_LIST_LINK; list <= OMP_LIST_LINK; list++)
- for (n = oc->clauses->lists[list]; n; n = n->next)
- {
- if (n->sym->mark)
- gfc_error ("Symbol %qs present on multiple clauses at %L",
- n->sym->name, &loc);
- else
- n->sym->mark = 1;
- }
- }
+ check_array_not_assumed (n->sym, n->where,
+ "DEVICE_RESIDENT");
+ }
- for (oc = ns->oacc_declare; oc; oc = oc->next)
- {
- for (list = OMP_LIST_LINK; list <= OMP_LIST_LINK; list++)
- for (n = oc->clauses->lists[list]; n; n = n->next)
- n->sym->mark = 0;
+ for (n = oc->clauses->lists[OMP_LIST_MAP]; n; n = n->next)
+ if (n->expr && n->expr->ref->type == REF_ARRAY)
+ gfc_error ("Subarray %qs is not allowed in $!ACC DECLARE at %L",
+ n->sym->name, &n->where);
}
+
+ /* Check for duplicate link, device_resident and data clauses. */
+ resolve_oacc_declare_map (ns->oacc_declare, OMP_LIST_LINK);
+ resolve_oacc_declare_map (ns->oacc_declare, OMP_LIST_DEVICE_RESIDENT);
+ resolve_oacc_declare_map (ns->oacc_declare, OMP_LIST_MAP);
}
@@ -5042,7 +4966,7 @@ gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED)
case EXEC_OACC_ENTER_DATA:
case EXEC_OACC_EXIT_DATA:
case EXEC_OACC_WAIT:
- resolve_omp_clauses (code, &code->loc, code->ext.omp_clauses, NULL,
+ resolve_omp_clauses (code, code->ext.omp_clauses, NULL,
true);
break;
case EXEC_OACC_PARALLEL_LOOP:
@@ -5104,11 +5028,11 @@ gfc_resolve_omp_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED)
case EXEC_OMP_TEAMS:
case EXEC_OMP_WORKSHARE:
if (code->ext.omp_clauses)
- resolve_omp_clauses (code, &code->loc, code->ext.omp_clauses, NULL);
+ resolve_omp_clauses (code, code->ext.omp_clauses, NULL);
break;
case EXEC_OMP_TARGET_UPDATE:
if (code->ext.omp_clauses)
- resolve_omp_clauses (code, &code->loc, code->ext.omp_clauses, NULL);
+ resolve_omp_clauses (code, code->ext.omp_clauses, NULL);
if (code->ext.omp_clauses == NULL
|| (code->ext.omp_clauses->lists[OMP_LIST_TO] == NULL
&& code->ext.omp_clauses->lists[OMP_LIST_FROM] == NULL))
@@ -5136,7 +5060,7 @@ gfc_resolve_omp_declare_simd (gfc_namespace *ns)
gfc_error ("!$OMP DECLARE SIMD should refer to containing procedure "
"%qs at %L", ns->proc_name->name, &ods->where);
if (ods->clauses)
- resolve_omp_clauses (NULL, &ods->where, ods->clauses, ns);
+ resolve_omp_clauses (NULL, ods->clauses, ns);
}
}
@@ -3634,12 +3634,65 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock,
return gfc_finish_block (&block);
}
-/* parallel loop and kernels loop. */
+/* Helper function to filter combined oacc constructs. ORIG_CLAUSES
+ contains the unfiltered list of clauses. LOOP_CLAUSES corresponds to
+ the filter list of loop clauses corresponding to the enclosed list.
+ This function is called recursively to handle device_type clauses. */
+
+static void
+gfc_filter_oacc_combined_clauses (gfc_omp_clauses **orig_clauses,
+ gfc_omp_clauses **loop_clauses)
+{
+ if (*orig_clauses == NULL)
+ {
+ *loop_clauses = NULL;
+ return;
+ }
+
+ *loop_clauses = gfc_get_omp_clauses ();
+
+ memset (*loop_clauses, 0, sizeof (gfc_omp_clauses));
+
+ (*loop_clauses)->gang = (*orig_clauses)->gang;
+ (*orig_clauses)->gang = false;
+ (*loop_clauses)->gang_expr = (*orig_clauses)->gang_expr;
+ (*orig_clauses)->gang_expr = NULL;
+ (*loop_clauses)->gang_static = (*orig_clauses)->gang_static;
+ (*orig_clauses)->gang_static = false;
+ (*loop_clauses)->vector = (*orig_clauses)->vector;
+ (*orig_clauses)->vector = false;
+ (*loop_clauses)->vector_expr = (*orig_clauses)->vector_expr;
+ (*orig_clauses)->vector_expr = NULL;
+ (*loop_clauses)->worker = (*orig_clauses)->worker;
+ (*orig_clauses)->worker = false;
+ (*loop_clauses)->worker_expr = (*orig_clauses)->worker_expr;
+ (*orig_clauses)->worker_expr = NULL;
+ (*loop_clauses)->seq = (*orig_clauses)->seq;
+ (*orig_clauses)->seq = false;
+ (*loop_clauses)->independent = (*orig_clauses)->independent;
+ (*orig_clauses)->independent = false;
+ (*loop_clauses)->par_auto = (*orig_clauses)->par_auto;
+ (*orig_clauses)->par_auto = false;
+ (*loop_clauses)->acc_collapse = (*orig_clauses)->acc_collapse;
+ (*orig_clauses)->acc_collapse = false;
+ (*loop_clauses)->collapse = (*orig_clauses)->collapse;
+ /* Don't reset (*orig_clauses)->collapse. */
+ (*loop_clauses)->tile = (*orig_clauses)->tile;
+ (*orig_clauses)->tile = false;
+ (*loop_clauses)->tile_list = (*orig_clauses)->tile_list;
+ (*orig_clauses)->tile_list = NULL;
+ (*loop_clauses)->device_types = (*orig_clauses)->device_types;
+
+ gfc_filter_oacc_combined_clauses (&(*orig_clauses)->dtype_clauses,
+ &(*loop_clauses)->dtype_clauses);
+}
+
+/* Combined OpenACC parallel loop and kernels loop. */
static tree
gfc_trans_oacc_combined_directive (gfc_code *code)
{
stmtblock_t block, inner, *pblock = NULL;
- gfc_omp_clauses construct_clauses, loop_clauses;
+ gfc_omp_clauses *loop_clauses;
tree stmt, oacc_clauses = NULL_TREE;
enum tree_code construct_code;
bool scan_nodesc_arrays = false;
@@ -3661,39 +3714,18 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
gfc_start_block (&block);
- memset (&loop_clauses, 0, sizeof (loop_clauses));
- if (code->ext.omp_clauses != NULL)
- {
- memcpy (&construct_clauses, code->ext.omp_clauses,
- sizeof (construct_clauses));
- loop_clauses.collapse = construct_clauses.collapse;
- loop_clauses.gang = construct_clauses.gang;
- loop_clauses.gang_expr = construct_clauses.gang_expr;
- loop_clauses.gang_static = construct_clauses.gang_static;
- loop_clauses.vector = construct_clauses.vector;
- loop_clauses.vector_expr = construct_clauses.vector_expr;
- loop_clauses.worker = construct_clauses.worker;
- loop_clauses.worker_expr = construct_clauses.worker_expr;
- loop_clauses.seq = construct_clauses.seq;
- loop_clauses.independent = construct_clauses.independent;
- construct_clauses.collapse = 0;
- construct_clauses.gang = false;
- construct_clauses.vector = false;
- construct_clauses.worker = false;
- construct_clauses.seq = false;
- construct_clauses.independent = false;
- oacc_clauses = gfc_trans_omp_clauses (&block, &construct_clauses,
- code->loc);
- }
+ gfc_filter_oacc_combined_clauses (&code->ext.omp_clauses, &loop_clauses);
+ oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
+ code->loc);
array_set = gfc_init_nodesc_arrays (&inner, &oacc_clauses, code,
scan_nodesc_arrays);
- if (!loop_clauses.seq)
+ if (!loop_clauses->seq)
pblock = (array_set && array_set->elements ()) ? &inner : █
else
pushlevel ();
- stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, &loop_clauses, NULL);
+ stmt = gfc_trans_omp_do (code, EXEC_OACC_LOOP, pblock, loop_clauses, NULL);
if (array_set && array_set->elements ())
gfc_add_expr_to_block (&inner, stmt);
@@ -3714,6 +3746,9 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
oacc_clauses);
gfc_add_expr_to_block (&block, stmt);
+ gfc_free_omp_clauses (loop_clauses);
+ code->ext.omp_clauses->device_types = NULL;
+
return gfc_finish_block (&block);
}
new file mode 100644
@@ -0,0 +1,322 @@
+/* Exercise *_parser_oacc_shape_clause by checking various combinations
+ of gang, worker and vector clause arguments. */
+
+/* { dg-compile } */
+
+int main ()
+{
+ int i;
+ int v = 32, w = 19;
+ int length = 1, num = 5;
+
+ /* Valid uses. */
+
+ #pragma acc kernels
+ #pragma acc loop gang worker vector
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(26)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length: 16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length: v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num: 16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num: v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(16)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(v)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: 16, num: 5)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: v, num: w)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num, static: 6)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: 5, num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(1, static:*)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static:*, 1)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(1, static:*)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: 5, static: 4)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: v, static: w)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num, static:num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length:length)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num:length)
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num:num)
+ for (i = 0; i < 10; i++)
+ ;
+
+ /* Invalid uses. */
+
+ #pragma acc kernels
+ #pragma acc loop gang(16, 24) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(v, w) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: 1, num:2, num:3, 4) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: 1 num:2, num:3, 4) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(1, num:2, num:3, 4) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num, num:5) /* { dg-error "unexpected argument" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(length:num) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(5, length:length) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(num:length) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(length:5) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(1, num:2) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static: * abc) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static:*num:1) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num: 5 static: *) /* { dg-error "expected '.' before" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(,static: *) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(,length:5) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(,num:10) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(,10) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(,10) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(,10) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(-12) /* { dg-warning "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(-1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num:-1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(num:1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static:-1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop gang(static:1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(-1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num:-1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop worker(num:1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(-1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length:-1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ #pragma acc kernels
+ #pragma acc loop vector(length:1.0) /* { dg-error "" } */
+ for (i = 0; i < 10; i++)
+ ;
+
+ return 0;
+}
@@ -6,7 +6,7 @@ template <typename T>
void f ()
{
extern T n ();
-#pragma omp parallel num_threads(n) // { dg-error "num_threads expression must be integral" }
+#pragma omp parallel num_threads(n) // { dg-error "'num_threads' expression must be integral" }
;
#pragma omp parallel for schedule(static, n)
for (int i = 0; i < 10; i++) // { dg-error "chunk size expression must be integral" }
@@ -6,7 +6,7 @@ template <typename T>
void f ()
{
T n = 6;
-#pragma omp parallel num_threads(n) // { dg-error "num_threads expression must be integral" }
+#pragma omp parallel num_threads(n) // { dg-error "'num_threads' expression must be integral" }
;
#pragma omp parallel for schedule(static, n)
for (int i = 0; i < 10; i++) // { dg-error "chunk size expression must be integral" }
new file mode 100644
@@ -0,0 +1,165 @@
+! Exercise combined OpenACC directives.
+
+! { dg-do compile }
+! { dg-options "-fopenacc -fdump-tree-omplower" }
+! { dg-prune-output "sorry, unimplemented: device_type clause is not supported yet" }
+
+subroutine test
+ implicit none
+ integer a(100), i, j, z
+
+ ! PARALLEL
+
+ !$acc parallel loop collapse (2)
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop gang
+ do i = 1, 100
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop worker
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop vector
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop seq
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop auto
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop tile (2, 3)
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop private (z) copy (a) gang device_type (nvidia) worker async (3) wait
+ do i = 1, 100
+ a(i) = i
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop independent
+ do i = 1, 100
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop private (z)
+ do i = 1, 100
+ z = 0
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel loop reduction (+:z) copy (z)
+ do i = 1, 100
+ end do
+ !$acc end parallel loop
+
+ ! KERNELS
+
+ !$acc kernels loop collapse (2)
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop gang
+ do i = 1, 100
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop worker
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop vector
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop seq
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop auto
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop tile (2, 3)
+ do i = 1, 100
+ do j = 1, 10
+ end do
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop private (z) copy (a) gang device_type (nvidia) worker async (3) wait
+ do i = 1, 100
+ a(i) = i
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop independent
+ do i = 1, 100
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop private (z)
+ do i = 1, 100
+ z = 0
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop reduction (+:z) copy (z)
+ do i = 1, 100
+ end do
+ !$acc end kernels loop
+end subroutine test
+
+! { dg-final { scan-tree-dump-times "acc loop collapse.2. private.j. private.i" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop gang" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop worker" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop vector" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop seq" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop auto" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "device_type.nvidia. . async.3. . map.force_tofrom:a" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop device_type.nvidia. . worker . gang private.i" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "private.z" 2 "omplower" } }
+! { dg-final { scan-tree-dump-times "map.force_tofrom:z .len: 4.. reduction..:z." 2 "omplower" } }
@@ -26,19 +26,29 @@ subroutine bsubr (foo)
integer, dimension (:) :: foo
- !$acc declare copy (foo) ! { dg-error "assumed-size dummy array" }
- !$acc declare copy (foo(1:2)) ! { dg-error "assumed-size dummy array" }
+ !$acc declare copy (foo) ! { dg-error "Assumed-size dummy array" }
+ !$acc declare copy (foo(1:2)) ! { dg-error "Assumed-size dummy array" }
-end subroutine
+end subroutine bsubr
-program test
- integer :: a(8)
+subroutine multiline
integer :: b(8)
+
+ !$acc declare copyin (b) ! { dg-error "present on multiple clauses" }
+ !$acc declare copyin (b)
+
+end subroutine multiline
+
+subroutine subarray
integer :: c(8)
+ !$acc declare copy (c(1:2)) ! { dg-error "Subarray 'c' is not allowed" }
+
+end subroutine subarray
+
+program test
+ integer :: a(8)
+
!$acc declare create (a) copyin (a) ! { dg-error "present on multiple clauses" }
- !$acc declare copyin (b)
- !$acc declare copyin (b) ! { dg-error "present on multiple clauses" }
- !$acc declare copy (c(1:2)) ! { dg-error "Subarray: 'c' not allowed" }
end program
new file mode 100644
@@ -0,0 +1,13 @@
+! Test if variable appearing in multiple clauses are errors.
+
+! { dg-compile }
+
+program combined
+ implicit none
+ integer a(100), i, j
+
+ !$acc parallel loop reduction (+:j) copy (j) copyout(j) ! { dg-error "Symbol 'j' present on multiple clauses" }
+ do i = 1, 100
+ end do
+ !$acc end parallel loop
+end program combined
@@ -11,6 +11,6 @@ subroutine foo (x)
!$omp simd linear (x) ! { dg-error "INTENT.IN. POINTER" }
do i = 1, 10
end do
-!$omp single ! { dg-error "INTENT.IN. POINTER" }
-!$omp end single copyprivate (x)
+!$omp single
+!$omp end single copyprivate (x) ! { dg-error "INTENT.IN. POINTER" }
end
deleted file mode 100644
@@ -1,37 +0,0 @@
-! { dg-do run }
-
-program main
- integer, parameter :: n = 32
- real :: a(n), b(n);
- integer :: i
-
- do i = 1, n
- a(i) = 1.0
- b(i) = 0.0
- end do
-
- !$acc parallel loop copy (a(1:n)) copy (b(1:n))
- do i = 1, n
- b(i) = 2.0
- a(i) = a(i) + b(i)
- end do
-
- do i = 1, n
- if (a(i) .ne. 3.0) call abort
-
- if (b(i) .ne. 2.0) call abort
- end do
-
- !$acc kernels loop copy (a(1:n)) copy (b(1:n))
- do i = 1, n
- b(i) = 3.0;
- a(i) = a(i) + b(i)
- end do
-
- do i = 1, n
- if (a(i) .ne. 6.0) call abort
-
- if (b(i) .ne. 3.0) call abort
- end do
-
-end program main
new file mode 100644
@@ -0,0 +1,39 @@
+! This test exercises combined directives.
+
+! { dg-do run }
+
+program main
+ integer, parameter :: n = 32
+ real :: a(n), b(n);
+ integer :: i
+
+ do i = 1, n
+ a(i) = 1.0
+ b(i) = 0.0
+ end do
+
+ !$acc parallel loop copy (a(1:n)) copy (b(1:n))
+ do i = 1, n
+ b(i) = 2.0
+ a(i) = a(i) + b(i)
+ end do
+
+ do i = 1, n
+ if (a(i) .ne. 3.0) call abort
+
+ if (b(i) .ne. 2.0) call abort
+ end do
+
+ !$acc kernels loop copy (a(1:n)) copy (b(1:n))
+ do i = 1, n
+ b(i) = 3.0;
+ a(i) = a(i) + b(i)
+ end do
+
+ do i = 1, n
+ if (a(i) .ne. 6.0) call abort
+
+ if (b(i) .ne. 3.0) call abort
+ end do
+
+end program main
This patch contains the following: * C front end changes from trunk: https://gcc.gnu.org/ml/gcc-patches/2015-10/msg02528.html * C++ front end changes from trunk: https://gcc.gnu.org/ml/gcc-patches/2015-10/msg02540.html * Proposed fortran cleanups and enhanced error reporting changes: https://gcc.gnu.org/ml/gcc-patches/2015-10/msg02288.html In addition, I've also added a couple of more test cases and updated the way that combined directives are handled in fortran. Because the device_type clauses form a chain of gfc_omp_clauses, I couldn't reuse gfc_split_omp_clauses for combined parallel and kernels loops. So that's why I introduced a new gfc_filter_oacc_combined_clauses function. I'll apply this patch to gomp-4_0-branch shortly. I know that I should have broken this patch down into smaller patches, but it was already arranged as one big patch in my source tree. Cesar