2015-12-16 Cesar Philippidis gcc/cp/ * cp-tree.h (bind_decls_match): Declare. * decl.c (bind_decls_match): New function. * parser.c (cp_parser_oacc_clause_bind): Remove TODOs and useless namespace checks. Defer string-ifying bind clauses with identifiers. Check for empty identifier strings. (cp_parser_oacc_routine): Clean up the pragma parser after detecting a non-pragma_external context. Remove useless namespace check. (cp_finalize_oacc_routine): Handle bind clauses with identifiers. gcc/testsuite/ * g++.dg/goacc/routine-2.C: Add more coverage. diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h index ea50048..cf3f022 100644 --- a/gcc/cp/cp-tree.h +++ b/gcc/cp/cp-tree.h @@ -5687,6 +5687,7 @@ extern void finish_scope (void); extern void push_switch (tree); extern void pop_switch (void); extern tree make_lambda_name (void); +extern int bind_decls_match (tree, tree); extern int decls_match (tree, tree); extern tree duplicate_decls (tree, tree, bool); extern tree declare_local_label (tree); diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c index e895c5a..3f8ecca 100644 --- a/gcc/cp/decl.c +++ b/gcc/cp/decl.c @@ -1121,6 +1121,138 @@ decls_match (tree newdecl, tree olddecl) return types_match; } +/* Similiar to decls_match, but only applies to FUNCTION_DECLS. Functions + in separate namespaces may match. +*/ + +int +bind_decls_match (tree newdecl, tree olddecl) +{ + int types_match; + + if (newdecl == olddecl) + return 1; + + if (TREE_CODE (newdecl) != TREE_CODE (olddecl)) + /* If the two DECLs are not even the same kind of thing, we're not + interested in their types. */ + return 0; + + gcc_assert (DECL_P (newdecl)); + gcc_assert (TREE_CODE (newdecl) == FUNCTION_DECL); + + tree f1 = TREE_TYPE (newdecl); + tree f2 = TREE_TYPE (olddecl); + tree p1 = TYPE_ARG_TYPES (f1); + tree p2 = TYPE_ARG_TYPES (f2); + tree r2; + + /* Specializations of different templates are different functions + even if they have the same type. */ + tree t1 = (DECL_USE_TEMPLATE (newdecl) + ? DECL_TI_TEMPLATE (newdecl) + : NULL_TREE); + tree t2 = (DECL_USE_TEMPLATE (olddecl) + ? DECL_TI_TEMPLATE (olddecl) + : NULL_TREE); + if (t1 != t2) + return 0; + + if (CP_DECL_CONTEXT (newdecl) != CP_DECL_CONTEXT (olddecl) + && TREE_CODE (CP_DECL_CONTEXT (newdecl)) != NAMESPACE_DECL + && TREE_CODE (CP_DECL_CONTEXT (olddecl)) != NAMESPACE_DECL + && ! (DECL_EXTERN_C_P (newdecl) + && DECL_EXTERN_C_P (olddecl))) + return 0; + + /* A new declaration doesn't match a built-in one unless it + is also extern "C". */ + if (DECL_IS_BUILTIN (olddecl) + && DECL_EXTERN_C_P (olddecl) && !DECL_EXTERN_C_P (newdecl)) + return 0; + + if (TREE_CODE (f1) != TREE_CODE (f2)) + return 0; + + /* A declaration with deduced return type should use its pre-deduction + type for declaration matching. */ + r2 = fndecl_declared_return_type (olddecl); + + if (same_type_p (TREE_TYPE (f1), r2)) + { + if (!prototype_p (f2) && DECL_EXTERN_C_P (olddecl) + && (DECL_BUILT_IN (olddecl) +#ifndef NO_IMPLICIT_EXTERN_C + || (DECL_IN_SYSTEM_HEADER (newdecl) && !DECL_CLASS_SCOPE_P (newdecl)) + || (DECL_IN_SYSTEM_HEADER (olddecl) && !DECL_CLASS_SCOPE_P (olddecl)) +#endif + )) + { + types_match = self_promoting_args_p (p1); + if (p1 == void_list_node) + TREE_TYPE (newdecl) = TREE_TYPE (olddecl); + } +#ifndef NO_IMPLICIT_EXTERN_C + else if (!prototype_p (f1) + && (DECL_EXTERN_C_P (olddecl) + && DECL_IN_SYSTEM_HEADER (olddecl) + && !DECL_CLASS_SCOPE_P (olddecl)) + && (DECL_EXTERN_C_P (newdecl) + && DECL_IN_SYSTEM_HEADER (newdecl) + && !DECL_CLASS_SCOPE_P (newdecl))) + { + types_match = self_promoting_args_p (p2); + TREE_TYPE (newdecl) = TREE_TYPE (olddecl); + } +#endif + else + types_match = + compparms (p1, p2) + && type_memfn_rqual (f1) == type_memfn_rqual (f2) + && (TYPE_ATTRIBUTES (TREE_TYPE (newdecl)) == NULL_TREE + || comp_type_attributes (TREE_TYPE (newdecl), + TREE_TYPE (olddecl)) != 0); + } + else + types_match = 0; + + /* The decls dont match if they correspond to two different versions + of the same function. Disallow extern "C" functions to be + versions for now. */ + if (types_match + && !DECL_EXTERN_C_P (newdecl) + && !DECL_EXTERN_C_P (olddecl) + && targetm.target_option.function_versions (newdecl, olddecl)) + { + /* Mark functions as versions if necessary. Modify the mangled decl + name if necessary. */ + if (DECL_FUNCTION_VERSIONED (newdecl) + && DECL_FUNCTION_VERSIONED (olddecl)) + return 0; + if (!DECL_FUNCTION_VERSIONED (newdecl)) + { + DECL_FUNCTION_VERSIONED (newdecl) = 1; + if (DECL_ASSEMBLER_NAME_SET_P (newdecl)) + mangle_decl (newdecl); + } + if (!DECL_FUNCTION_VERSIONED (olddecl)) + { + DECL_FUNCTION_VERSIONED (olddecl) = 1; + if (DECL_ASSEMBLER_NAME_SET_P (olddecl)) + mangle_decl (olddecl); + } + cgraph_node::record_function_versions (olddecl, newdecl); + return 0; + } + + // Normal functions can be constrained, as can variable partial + // specializations. + if (types_match && VAR_OR_FUNCTION_DECL_P (newdecl)) + types_match = equivalently_constrained (newdecl, olddecl); + + return types_match; +} + /* If NEWDECL is `static' and an `extern' was seen previously, warn about it. OLDDECL is the previous declaration. diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 6556db3..824261e 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -31552,7 +31552,6 @@ cp_parser_oacc_clause_bind (cp_parser *parser, tree list) cp_token *token = cp_lexer_peek_token (parser->lexer); if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)) { - //TODO tree id = cp_parser_id_expression (parser, /*template_p=*/false, /*check_dependency_p=*/true, /*template_p=*/NULL, @@ -31562,44 +31561,40 @@ cp_parser_oacc_clause_bind (cp_parser *parser, tree list) if (id != error_mark_node && decl == error_mark_node) cp_parser_name_lookup_error (parser, id, decl, NLE_NULL, token->location); - if (/* TODO */ !decl || decl == error_mark_node) + if (!decl || decl == error_mark_node) error_at (token->location, "%qE has not been declared", token->u.value); - else if (/* TODO */ is_overloaded_fn (decl) + else if (is_overloaded_fn (decl) && (TREE_CODE (decl) != FUNCTION_DECL || DECL_FUNCTION_TEMPLATE_P (decl))) error_at (token->location, "%qE names a set of overloads", token->u.value); - else if (/* TODO */ !DECL_NAMESPACE_SCOPE_P (decl)) - { - /* Perhaps we should use the same rule as declarations in different - namespaces? */ - error_at (token->location, - "%qE does not refer to a namespace scope function", - token->u.value); - } else if (TREE_CODE (decl) != FUNCTION_DECL) error_at (token->location, "%qE does not refer to a function", token->u.value); else - { - //TODO? TREE_USED (decl) = 1; - tree name_id = DECL_NAME (decl); - name = build_string (IDENTIFIER_LENGTH (name_id), - IDENTIFIER_POINTER (name_id)); - } - //cp_lexer_consume_token (parser->lexer); + name = decl; } else if (cp_lexer_next_token_is (parser->lexer, CPP_STRING)) { name = token->u.value; cp_lexer_consume_token (parser->lexer); + + /* This shouldn't be an empty string. */ + if (strcmp (TREE_STRING_POINTER (name), "\"\"") == 0) + error_at (token->location, + "bind argument must not be an empty string"); + + parser->translate_strings_p = save_translate_strings_p; } else - cp_parser_error (parser, - "expected identifier or character string literal"); - parser->translate_strings_p = save_translate_strings_p; + { + cp_parser_error (parser, + "expected identifier or character string literal"); + cp_parser_skip_to_closing_parenthesis (parser, false, false, true); + return list; + } cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN); if (name != error_mark_node) { @@ -36059,6 +36054,7 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, if (context != pragma_external) { cp_parser_error (parser, "%<#pragma acc routine%> not at file scope"); + cp_parser_skip_to_pragma_eol (parser, pragma_tok); parser->oacc_routine->error_seen = true; parser->oacc_routine = NULL; return; @@ -36130,16 +36126,6 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok, return; } - /* Perhaps we should use the same rule as declarations in different - namespaces? */ - if (!DECL_NAMESPACE_SCOPE_P (decl)) - { - error_at (loc, "%<#pragma acc routine%> does not refer to a " - "namespace scope function"); - parser->oacc_routine = NULL; - return; - } - if (!decl || TREE_CODE (decl) != FUNCTION_DECL) { error_at (loc, @@ -36288,6 +36274,34 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn) parser->oacc_routine = NULL; } + /* Process the bind clause, if present. */ + for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_BIND) + continue; + + tree bind_decl = OMP_CLAUSE_BIND_NAME (c); + + /* String arguments don't require any special treatment. */ + if (TREE_CODE (bind_decl) != FUNCTION_DECL) + break; + + if (!bind_decls_match (bind_decl, fndecl)) + { + error_at (OMP_CLAUSE_LOCATION (c), + "bind identifier %qE is not compatible with " + "function %qE", bind_decl, fndecl); + return; + } + + tree name_id = decl_assembler_name (bind_decl); + tree name = build_string (IDENTIFIER_LENGTH (name_id), + IDENTIFIER_POINTER (name_id)); + OMP_CLAUSE_BIND_NAME (c) = name; + + break; + } + /* Process for function attrib */ tree dims = build_oacc_routine_dims (clauses); replace_oacc_fn_attrib (fndecl, dims); diff --git a/gcc/testsuite/g++.dg/goacc/routine-2.C b/gcc/testsuite/g++.dg/goacc/routine-2.C index 92fc161..44d404a 100644 --- a/gcc/testsuite/g++.dg/goacc/routine-2.C +++ b/gcc/testsuite/g++.dg/goacc/routine-2.C @@ -14,3 +14,123 @@ one() return 1; } #pragma acc routine (one) bind(one_d) /* { dg-error "names a set of overloads" } */ + +int incr (int); +float incr (float); +int inc; + +#pragma acc routine (incr) /* { dg-error "names a set of overloads" } */ + +#pragma acc routine (increment) /* { dg-error "has not been declared" } */ + +#pragma acc routine (inc) /* { dg-error "does not refer to a function" } */ + +#pragma acc routine (+) /* { dg-error "expected unqualified-id before '.' token" } */ + +int sum (int, int); + +namespace foo { +#pragma acc routine (sum) + int sub (int, int); +} + +#pragma acc routine (foo::sub) + +/* It's strange to apply a routine directive to subset of overloaded + functions, but that is permissible in OpenACC 2.x. */ + +int decr (int a); + +#pragma acc routine +float decr (float a); + +/* Bind clause. */ + +#pragma acc routine +float +mult (float a, float b) +{ + return a * b; +} + +#pragma acc routine bind(mult) /* { dg-error "bind identifier 'mult' is not compatible with function 'broken_mult1'" } */ +float +broken_mult1 (int a, int b) +{ + return a + b; +} + +/* This should result in a link error, but it's valid for a compile test. */ +#pragma acc routine bind("mult") +float +broken_mult2 (float a, float b) +{ + return a + b; +} + +#pragma acc routine bind(sum2) /* { dg-error "'sum2' has not been declared" } */ +int broken_mult3 (int a, int b); + +#pragma acc routine bind(foo::sub) +int broken_mult4 (int a, int b); + +namespace bar +{ +#pragma acc routine bind (mult) + float working_mult (float a, float b) + { + return a * b; + } +} + +#pragma acc routine +int div (int, int); + +#pragma acc routine +float div (float, float); + +#pragma acc routine bind (div) /* { dg-error "'div' names a set of overloads" } */ +float +my_div (float a, float b) +{ + return a / b; +} + +#pragma acc routine bind (other_div) /* { dg-error "'other_div' has not been declared" } */ +float +my_div2 (float a, float b) +{ + return a / b; +} + +int div_var; + +#pragma acc routine bind (div_var) /* { dg-error "'div_var' does not refer to a function" } */ +float my_div3 (float, float); + +#pragma acc routine bind (div_var) /* { dg-error "'div_var' does not refer to a function" } */ +float my_div4 (float, float); + +#pragma acc routine bind (%) /* { dg-error "expected identifier or character string literal" } */ +float my_div5 (float, float); + +#pragma acc routine bind ("") /* { dg-error "bind argument must not be an empty string" } */ +float my_div6 (float, float); + +struct astruct +{ + #pragma acc routine /* { dg-error "not at file scope before end of line" } */ + int sum (int a, int b) + { + return a + b; + } +}; + +class aclass +{ + #pragma acc routine /* { dg-error "not at file scope before end of line" } */ + int sum (int a, int b) + { + return a + b; + } +};