* [gomp4] fortran cleanups and c/c++ loop parsing backport
@ 2015-10-27 18:54 Cesar Philippidis
2015-10-28 11:03 ` Thomas Schwinge
2015-11-04 11:39 ` [gomp4] fortran cleanups and c/c++ loop parsing backport Thomas Schwinge
0 siblings, 2 replies; 7+ messages in thread
From: Cesar Philippidis @ 2015-10-27 18:54 UTC (permalink / raw)
To: gcc-patches; +Cc: james norris
[-- Attachment #1: Type: text/plain, Size: 904 bytes --]
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
[-- Attachment #2: gfc_combined_directives_plus_backport.diff --]
[-- Type: text/x-patch, Size: 70068 bytes --]
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.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 3c36fc6..a1465bf 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -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");
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 3db6d0a..d113cfb 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -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");
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index a6ee58a..9db8b27 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -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)
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 312e30d..5d58d2b 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -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;
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index c42a2c2..9621eaf 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -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);
}
}
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index de2422f..bec2de4 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -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);
}
diff --git a/gcc/testsuite/c-c++-common/goacc/loop-shape.c b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
new file mode 100644
index 0000000..b6d3156
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/loop-shape.c
@@ -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;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/pr33372-1.C b/gcc/testsuite/g++.dg/gomp/pr33372-1.C
index 62900bf..e9da259 100644
--- a/gcc/testsuite/g++.dg/gomp/pr33372-1.C
+++ b/gcc/testsuite/g++.dg/gomp/pr33372-1.C
@@ -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" }
diff --git a/gcc/testsuite/g++.dg/gomp/pr33372-3.C b/gcc/testsuite/g++.dg/gomp/pr33372-3.C
index 8220f3c..f0a1910 100644
--- a/gcc/testsuite/g++.dg/gomp/pr33372-3.C
+++ b/gcc/testsuite/g++.dg/gomp/pr33372-3.C
@@ -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" }
diff --git a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
new file mode 100644
index 0000000..a72090c
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
@@ -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" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95
index afdbe2e..f9ffe9e 100644
--- a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95
@@ -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
diff --git a/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90 b/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90
new file mode 100644
index 0000000..2870076
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/multi-clause.f90
@@ -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
diff --git a/gcc/testsuite/gfortran.dg/gomp/intentin1.f90 b/gcc/testsuite/gfortran.dg/gomp/intentin1.f90
index f2a2e98..8bd53aa 100644
--- a/gcc/testsuite/gfortran.dg/gomp/intentin1.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/intentin1.f90
@@ -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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
deleted file mode 100644
index 0cd8a67..0000000
--- a/libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
+++ /dev/null
@@ -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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/combined-directive-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/combined-directive-1.f90
new file mode 100644
index 0000000..94100b2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/combined-directive-1.f90
@@ -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
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [gomp4] fortran cleanups and c/c++ loop parsing backport
2015-10-27 18:54 [gomp4] fortran cleanups and c/c++ loop parsing backport Cesar Philippidis
@ 2015-10-28 11:03 ` Thomas Schwinge
2015-10-28 15:35 ` Cesar Philippidis
2015-11-04 11:39 ` [gomp4] fortran cleanups and c/c++ loop parsing backport Thomas Schwinge
1 sibling, 1 reply; 7+ messages in thread
From: Thomas Schwinge @ 2015-10-28 11:03 UTC (permalink / raw)
To: Cesar Philippidis, gcc-patches; +Cc: james norris
[-- Attachment #1: Type: text/plain, Size: 11239 bytes --]
Hi Cesar!
On Tue, 27 Oct 2015 11:36:10 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> 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
I suppose the latter is a prerequisite for other Fortran front end
changes you've also committed? Otherwise, why not get that patch into
trunk first? That sould save me from having to deal with potentially
more merge conflicts later on...
> 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.
Thanks, but...
> I'll apply this patch to gomp-4_0-branch shortly. I know that I should
> have broken this patch down into smaller patches
Yes.
> but it was already
> arranged as one big patch in my source tree.
You're using Git, so that's not a good excuse. :-P
> --- a/gcc/fortran/trans-openmp.c
> +++ b/gcc/fortran/trans-openmp.c
> @@ -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));
This has already been created zero-initialized.
> + (*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. */
Why? (Extend source code comment?) The original code (cited just below)
did this differently.
> + (*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);
>
With...
> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
... newly added, and...
> --- a/libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
> +++ /dev/null
... renamed to...
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/combined-directive-1.f90
... this, plus the unaltered
libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c, we now got three
variants: "combined-directives", "combined-directive", and "combdir".
Please settle on one of them.
> --- a/gcc/testsuite/gfortran.dg/gomp/intentin1.f90
> +++ b/gcc/testsuite/gfortran.dg/gomp/intentin1.f90
> @@ -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
Please revert this unrelated change.
Additionally, in r229482 I checked in the following to restore bootstrap
builds:
[...]/source-gcc/gcc/cp/parser.c:29649:1: error: 'tree_node* require_positive_expr(tree, location_t, const char*)' defined but not used [-Werror=unused-function]
require_positive_expr (tree t, location_t loc, const char *str)
^
[...]/source-gcc/gcc/fortran/openmp.c: In function 'match gfc_match_oacc_declare()':
[...]/source-gcc/gcc/fortran/openmp.c:1449:30: error: unused variable 'oc' [-Werror=unused-variable]
gfc_oacc_declare *new_oc, *oc;
^
[...]/source-gcc/gcc/fortran/openmp.c: In function 'void gfc_resolve_oacc_declare(gfc_namespace*)':
[...]/source-gcc/gcc/fortran/openmp.c:4918:7: error: unused variable 'list' [-Werror=unused-variable]
int list;
^
commit 1e53d795ea68bcc7e5d5a7fa21e58c506e557cb4
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Oct 28 10:57:45 2015 +0000
Address -Werror=unused-variable and -Werror=unused-function diagnostics
gcc/cp/
* parser.c (require_positive_expr): Remove function.
gcc/fortran/
* openmp.c (gfc_match_oacc_declare): Remove oc local variable.
(gfc_resolve_oacc_declare): Remove list local variable.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@229482 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/cp/ChangeLog.gomp | 4 ++++
gcc/cp/parser.c | 17 -----------------
gcc/fortran/ChangeLog.gomp | 5 +++++
gcc/fortran/openmp.c | 3 +--
4 files changed, 10 insertions(+), 19 deletions(-)
diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp
index c0b96ba..b96bfce 100644
--- gcc/cp/ChangeLog.gomp
+++ gcc/cp/ChangeLog.gomp
@@ -1,3 +1,7 @@
+2015-10-28 Thomas Schwinge <thomas@codesourcery.com>
+
+ * parser.c (require_positive_expr): Remove function.
+
2015-10-27 Cesar Philippidis <cesar@codesourcery.com>
* parser.c (cp_parser_oacc_shape_clause): Backport from trunk.
diff --git gcc/cp/parser.c gcc/cp/parser.c
index d113cfb..2c6060b 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -29642,23 +29642,6 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list)
return list;
}
-/* Attempt to statically determine when the number T isn't positive.
- Warn if we determined this and return positive one as the new
- expression. */
-static tree
-require_positive_expr (tree t, location_t loc, const char *str)
-{
- tree c = fold_build2_loc (loc, LE_EXPR, boolean_type_node, t,
- build_int_cst (TREE_TYPE (t), 0));
- if (c == boolean_true_node)
- {
- warning_at (loc, 0,
- "%<%s%> value must be positive", str);
- t = integer_one_node;
- }
- return t;
-}
-
/* OpenACC:
num_gangs ( expression )
num_workers ( expression )
diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp
index d126367..9e1b95b 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2015-10-28 Thomas Schwinge <thomas@codesourcery.com>
+
+ * openmp.c (gfc_match_oacc_declare): Remove oc local variable.
+ (gfc_resolve_oacc_declare): Remove list local variable.
+
2015-10-27 Cesar Philippidis <cesar@codesourcery.com>
* gfortran.h (gfc_omp_namelist): Add locus where member.
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index 9621eaf..afcce9a 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -1446,7 +1446,7 @@ gfc_match_oacc_declare (void)
gfc_omp_clauses *c;
gfc_omp_namelist *n;
gfc_namespace *ns = gfc_current_ns;
- gfc_oacc_declare *new_oc, *oc;
+ gfc_oacc_declare *new_oc;
bool module_var = false;
if (gfc_match_omp_clauses (&c, OACC_DECLARE_CLAUSES, 0, false, false, true)
@@ -4915,7 +4915,6 @@ resolve_oacc_declare_map (gfc_oacc_declare *declare, int list)
void
gfc_resolve_oacc_declare (gfc_namespace *ns)
{
- int list;
gfc_omp_namelist *n;
locus loc;
gfc_oacc_declare *oc;
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [gomp4] fortran cleanups and c/c++ loop parsing backport
2015-10-28 11:03 ` Thomas Schwinge
@ 2015-10-28 15:35 ` Cesar Philippidis
2015-10-29 9:06 ` Thomas Schwinge
0 siblings, 1 reply; 7+ messages in thread
From: Cesar Philippidis @ 2015-10-28 15:35 UTC (permalink / raw)
To: Thomas Schwinge, gcc-patches; +Cc: james norris
[-- Attachment #1: Type: text/plain, Size: 8542 bytes --]
On 10/28/2015 04:00 AM, Thomas Schwinge wrote:
> Hi Cesar!
>
> On Tue, 27 Oct 2015 11:36:10 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
>> 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
>
> I suppose the latter is a prerequisite for other Fortran front end
> changes you've also committed? Otherwise, why not get that patch into
> trunk first? That sould save me from having to deal with potentially
> more merge conflicts later on...
It wasn't necessarily a prerequisite for these changes, but I've been
trying to get that patch into trunk for a while now. Plus, part of those
cleanups touched declare, which Jim is going to work on soon.
>> 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.
>
> Thanks, but...
>
>> I'll apply this patch to gomp-4_0-branch shortly. I know that I should
>> have broken this patch down into smaller patches
>
> Yes.
>
>> but it was already
>> arranged as one big patch in my source tree.
>
> You're using Git, so that's not a good excuse. :-P
I find git to be too temperamental.
>> --- a/gcc/fortran/trans-openmp.c
>> +++ b/gcc/fortran/trans-openmp.c
>> @@ -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));
>
> This has already been created zero-initialized.
I was just doing what I was doing before. I removed that in the follow
up patch.
>> + (*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. */
>
> Why? (Extend source code comment?) The original code (cited just below)
> did this differently.
Because that's what gfc_split_omp_clauses does. I'm not sure what that's
required for gfc_trans_omp_do, but it is. gfc_trans_omp_do appears to be
operating on two sets of clauses for some non-obvious reason.
>> + (*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);
>>
>
>
> With...
>
>> --- /dev/null
>> +++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
>
> ... newly added, and...
>
>> --- a/libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
>> +++ /dev/null
>
> ... renamed to...
>
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.oacc-fortran/combined-directive-1.f90
>
> ... this, plus the unaltered
> libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c, we now got three
> variants: "combined-directives", "combined-directive", and "combdir".
> Please settle on one of them.
The problem here was I didn't know what compdir stood for, so I renamed
it to combined-directive-foo. I guess I didn't keep the name consistent
when I introduced a new compiler time test. The attached patch adjusts
the test in libgomp to make it consistent with it's compile time test.
>> --- a/gcc/testsuite/gfortran.dg/gomp/intentin1.f90
>> +++ b/gcc/testsuite/gfortran.dg/gomp/intentin1.f90
>> @@ -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
>
> Please revert this unrelated change.
This is related to the fortran patch I referenced above. That fortran
patch did two things:
1. It introduced a function to detect if variables appear in in
omp map lists. (Because we don't support 'acc copyin (foo)
copyout (foo)')
2. It associates a gfc_locus with each entry inside an omp map list.
That's why I had to adjust that test case.
I could revert that patch from gomp4, but it really impacts declare.
There's so much copy-and-paste going on there. What do you want to do
with this?
Cesar
[-- Attachment #2: gfc_combined_directives_plus_backport-followup.diff --]
[-- Type: text/x-patch, Size: 3420 bytes --]
2015-10-28 Cesar Philippidis <cesar@codesourcery.com>
gcc/fortran/
* trans-openmp.c (gfc_filter_oacc_combined_clauses): Don't zero-
initialize loop_clauses when it's already initialized.
libgomp/
* testsuite/libgomp.oacc-fortran/combined-directive-1.f90: Rename to...
* testsuite/libgomp.oacc-fortran/combined-directives-1.f90: ... this.
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index bec2de4..a4466ed 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -3651,8 +3651,6 @@ gfc_filter_oacc_combined_clauses (gfc_omp_clauses **orig_clauses,
*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;
@@ -3676,7 +3674,9 @@ gfc_filter_oacc_combined_clauses (gfc_omp_clauses **orig_clauses,
(*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. */
+ /* Don't reset (*orig_clauses)->collapse. It should be present on
+ both the kernels/parallel and loop constructs, similar to how
+ gfc_split_omp_clauses duplicates it in combined OMP constructs. */
(*loop_clauses)->tile = (*orig_clauses)->tile;
(*orig_clauses)->tile = false;
(*loop_clauses)->tile_list = (*orig_clauses)->tile_list;
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/combined-directive-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/combined-directive-1.f90
deleted file mode 100644
index 94100b2..0000000
--- a/libgomp/testsuite/libgomp.oacc-fortran/combined-directive-1.f90
+++ /dev/null
@@ -1,39 +0,0 @@
-! 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
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/combined-directives-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/combined-directives-1.f90
new file mode 100644
index 0000000..94100b2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/combined-directives-1.f90
@@ -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
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [gomp4] fortran cleanups and c/c++ loop parsing backport
2015-10-28 15:35 ` Cesar Philippidis
@ 2015-10-29 9:06 ` Thomas Schwinge
2015-10-29 9:25 ` Improve filenames for test cases of OpenACC combined directives (was: [gomp4] fortran cleanups and c/c++ loop parsing backport) Thomas Schwinge
0 siblings, 1 reply; 7+ messages in thread
From: Thomas Schwinge @ 2015-10-29 9:06 UTC (permalink / raw)
To: Cesar Philippidis, Jakub Jelinek; +Cc: james norris, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 6012 bytes --]
Hi!
On Wed, 28 Oct 2015 08:30:56 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> On 10/28/2015 04:00 AM, Thomas Schwinge wrote:
> > On Tue, 27 Oct 2015 11:36:10 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> >> 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
> >
> > I suppose the latter is a prerequisite for other Fortran front end
> > changes you've also committed? Otherwise, why not get that patch into
> > trunk first? That sould save me from having to deal with potentially
> > more merge conflicts later on...
>
> It wasn't necessarily a prerequisite for these changes, but I've been
> trying to get that patch into trunk for a while now.
Generally, please get such "widespread" patches (touching a lot of code,
even if trivially) into trunk first. With every merge from trunk into
gomp-4_0-branch I'll now have to make sure to adjust every new trunk code
to cooperate with your patch.
Jakub, can you please review
<http://news.gmane.org/find-root.php?message_id=%3C5628FEFF.50809%40codesourcery.com%3E>?
> Plus, part of those
> cleanups touched declare, which Jim is going to work on soon.
I don't understand how that relates.
Anyway, the patch is now committed on gomp-4_0-branch, and hopefully will
land in trunk soon; let's move on.
> >> --- a/gcc/fortran/trans-openmp.c
> >> +++ b/gcc/fortran/trans-openmp.c
> >> +/* 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)
> >> + (*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. */
> >
> > Why? (Extend source code comment?) The original code (cited just below)
> > did this differently.
>
> Because that's what gfc_split_omp_clauses does. I'm not sure what that's
> required for gfc_trans_omp_do, but it is. gfc_trans_omp_do appears to be
> operating on two sets of clauses for some non-obvious reason.
>
> >> @@ -3661,39 +3714,18 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
> >> - loop_clauses.collapse = construct_clauses.collapse;
> >> - [...]
> >> - construct_clauses.collapse = 0;
(The "old" behavior is also what the C/C++ front ends are doing, by the
way, so that would also need to be updated.)
Assuming I'm looking at the correct thing, the collapse clause handling
in gfc_split_omp_clauses says "duplicate collapse" which is more
expressive to me than "don't reset (*orig_clauses)->collapse". I haven't
researched the requirement for OpenMP, but I don't see how a collapse
clause would be relevant to have on an OpenACC parallel or kernels
construct? Which detail do I overlook?
> > With...
> >
> >> --- /dev/null
> >> +++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
> >
> > ... newly added, and...
> >
> >> --- a/libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
> >> +++ /dev/null
> >
> > ... renamed to...
> >
> >> --- /dev/null
> >> +++ b/libgomp/testsuite/libgomp.oacc-fortran/combined-directive-1.f90
> >
> > ... this, plus the unaltered
> > libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c, we now got three
> > variants: "combined-directives", "combined-directive", and "combdir".
> > Please settle on one of them.
>
> The problem here was I didn't know what compdir stood for
*Comb*ined *dir*ectives, I guessed.
> so I renamed
> it to combined-directive-foo. I guess I didn't keep the name consistent
> when I introduced a new compiler time test. The attached patch adjusts
> the test in libgomp to make it consistent with it's compile time test.
Again, to not make merges needlessly difficult, the preference is to do
such changes on trunk first. (Which I'll now do in the following, also
taking care to rename
libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c in the same
way...).
> >> --- a/gcc/testsuite/gfortran.dg/gomp/intentin1.f90
> >> +++ b/gcc/testsuite/gfortran.dg/gomp/intentin1.f90
> >> @@ -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
> >
> > Please revert this unrelated change.
>
> This is related to the fortran patch I referenced above. That fortran
> patch did two things:
>
> 1. It introduced a function to detect if variables appear in in
> omp map lists. (Because we don't support 'acc copyin (foo)
> copyout (foo)')
>
> 2. It associates a gfc_locus with each entry inside an omp map list.
> That's why I had to adjust that test case.
>
> I could revert that patch from gomp4, but it really impacts declare.
> There's so much copy-and-paste going on there. What do you want to do
> with this?
Copy-and-paste?
Anyway, this is my fault: it looked to me like just whitespace change,
and I didn't realize you'd moved the dg-error directive from one line to
another; all good then.
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 7+ messages in thread
* Improve filenames for test cases of OpenACC combined directives (was: [gomp4] fortran cleanups and c/c++ loop parsing backport)
2015-10-29 9:06 ` Thomas Schwinge
@ 2015-10-29 9:25 ` Thomas Schwinge
0 siblings, 0 replies; 7+ messages in thread
From: Thomas Schwinge @ 2015-10-29 9:25 UTC (permalink / raw)
To: gcc-patches; +Cc: Cesar Philippidis, Jakub Jelinek, james norris
[-- Attachment #1: Type: text/plain, Size: 4573 bytes --]
Hi!
On Thu, 29 Oct 2015 09:56:18 +0100, I wrote:
> On Wed, 28 Oct 2015 08:30:56 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> > On 10/28/2015 04:00 AM, Thomas Schwinge wrote:
> > > On Tue, 27 Oct 2015 11:36:10 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> > > With...
> > >
> > >> --- /dev/null
> > >> +++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
> > >
> > > ... newly added, and...
> > >
> > >> --- a/libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
> > >> +++ /dev/null
> > >
> > > ... renamed to...
> > >
> > >> --- /dev/null
> > >> +++ b/libgomp/testsuite/libgomp.oacc-fortran/combined-directive-1.f90
> > >
> > > ... this, plus the unaltered
> > > libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c, we now got three
> > > variants: "combined-directives", "combined-directive", and "combdir".
> > > Please settle on one of them.
> >
> > The problem here was I didn't know what compdir stood for
>
> *Comb*ined *dir*ectives, I guessed.
>
> > so I renamed
> > it to combined-directive-foo. I guess I didn't keep the name consistent
> > when I introduced a new compiler time test. The attached patch adjusts
> > the test in libgomp to make it consistent with it's compile time test.
>
> Again, to not make merges needlessly difficult, the preference is to do
> such changes on trunk first. (Which I'll now do in the following, also
> taking care to rename
> libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c in the same
> way...).
As obvious, committed to trunk in r229518:
commit e8ec2b06072a532b8d7319ed598ac1dad2584fea
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Oct 29 09:03:40 2015 +0000
Improve filenames for test cases of OpenACC combined directives
libgomp/
* testsuite/libgomp.oacc-c-c++-common/combdir-1.c: Rename to...
* testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c:
... this. Add a description of the test at the top of the file.
* testsuite/libgomp.oacc-fortran/combdir-1.f90: Rename file to...
* testsuite/libgomp.oacc-fortran/combined-directives-1.f90:
... this. Add a description of the test at the top of the file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@229518 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 10 ++++++++++
.../{combdir-1.c => combined-directives-1.c} | 2 ++
.../{combdir-1.f90 => combined-directives-1.f90} | 2 ++
3 files changed, 14 insertions(+)
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 8f44af0..c78881b 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,3 +1,13 @@
+2015-10-29 Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/combdir-1.c: Rename to...
+ * testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c:
+ ... this. Add a description of the test at the top of the file.
+ * testsuite/libgomp.oacc-fortran/combdir-1.f90: Rename file to...
+ * testsuite/libgomp.oacc-fortran/combined-directives-1.f90:
+ ... this. Add a description of the test at the top of the file.
+
2015-10-28 Nathan Sidwell <nathan@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/loop-g-1.c: New.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c
similarity index 93%
rename from libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c
index a7def92..dad6d13 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/combdir-1.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/combined-directives-1.c
@@ -1,3 +1,5 @@
+/* This test exercises combined directives. */
+
/* { dg-do run } */
#include <stdlib.h>
diff --git libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90 libgomp/testsuite/libgomp.oacc-fortran/combined-directives-1.f90
similarity index 92%
rename from libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
rename to libgomp/testsuite/libgomp.oacc-fortran/combined-directives-1.f90
index 0cd8a67..94100b2 100644
--- libgomp/testsuite/libgomp.oacc-fortran/combdir-1.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/combined-directives-1.f90
@@ -1,3 +1,5 @@
+! This test exercises combined directives.
+
! { dg-do run }
program main
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [gomp4] fortran cleanups and c/c++ loop parsing backport
2015-10-27 18:54 [gomp4] fortran cleanups and c/c++ loop parsing backport Cesar Philippidis
2015-10-28 11:03 ` Thomas Schwinge
@ 2015-11-04 11:39 ` Thomas Schwinge
2015-11-04 14:49 ` Cesar Philippidis
1 sibling, 1 reply; 7+ messages in thread
From: Thomas Schwinge @ 2015-11-04 11:39 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: gcc-patches, james norris
[-- Attachment #1: Type: text/plain, Size: 7135 bytes --]
Hi Cesar!
On Tue, 27 Oct 2015 11:36:10 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> * Proposed fortran cleanups and enhanced error reporting changes:
> https://gcc.gnu.org/ml/gcc-patches/2015-10/msg02288.html
... has now been applied to trunk, in an altered version, so we now got
some divergence between trunk and gomp-4_0-branch to resolve.
> 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.
> --- a/gcc/fortran/openmp.c
> +++ b/gcc/fortran/openmp.c
> @@ -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;
> }
Shouldn't the "%C" -> "%L (&n->where)" substitution also be done for that
one? If yes, please do so; if no, please add a source code comment, why.
>
> 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;
> }
> -/* 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)
If I understand correctly, that has not been approved for trunk, so we
should revert this on gomp-4_0-branch, too?
> {
> 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;
> + }
> }
> @@ -4966,64 +4927,27 @@ gfc_resolve_oacc_declare (gfc_namespace *ns)
> {
> loc = oc->where;
As far as I can tell, given the following changes, the "loc" local
variable is now unused, so should be removed?
>
> - 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);
> }
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [gomp4] fortran cleanups and c/c++ loop parsing backport
2015-11-04 11:39 ` [gomp4] fortran cleanups and c/c++ loop parsing backport Thomas Schwinge
@ 2015-11-04 14:49 ` Cesar Philippidis
0 siblings, 0 replies; 7+ messages in thread
From: Cesar Philippidis @ 2015-11-04 14:49 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches, james norris
On 11/04/2015 03:39 AM, Thomas Schwinge wrote:
> On Tue, 27 Oct 2015 11:36:10 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
>> * Proposed fortran cleanups and enhanced error reporting changes:
>> https://gcc.gnu.org/ml/gcc-patches/2015-10/msg02288.html
>
> ... has now been applied to trunk, in an altered version, so we now got
> some divergence between trunk and gomp-4_0-branch to resolve.
I've been concentrating on trunk lately. I'll backport these changes
when my other fortran changes go in
(https://gcc.gnu.org/ml/gcc-patches/2015-11/msg00287.html).
Cesar
^ permalink raw reply [flat|nested] 7+ messages in thread
end of thread, other threads:[~2015-11-04 14:49 UTC | newest]
Thread overview: 7+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-10-27 18:54 [gomp4] fortran cleanups and c/c++ loop parsing backport Cesar Philippidis
2015-10-28 11:03 ` Thomas Schwinge
2015-10-28 15:35 ` Cesar Philippidis
2015-10-29 9:06 ` Thomas Schwinge
2015-10-29 9:25 ` Improve filenames for test cases of OpenACC combined directives (was: [gomp4] fortran cleanups and c/c++ loop parsing backport) Thomas Schwinge
2015-11-04 11:39 ` [gomp4] fortran cleanups and c/c++ loop parsing backport Thomas Schwinge
2015-11-04 14:49 ` Cesar Philippidis
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).