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
2015-10-27 Cesar Philippidis <ce...@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