2015-12-16 Cesar Philippidis <cesar@codesourcery.com>
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.
@@ -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);
@@ -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.
@@ -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);
@@ -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;
+ }
+};