On 12/14/2015 12:36 PM, Cesar Philippidis wrote:
> On 12/08/2015 11:55 AM, Thomas Schwinge wrote:
>> On Sat, 14 Nov 2015 09:36:36 +0100, I wrote:
>> C front end:
>>
>> --- gcc/c/c-parser.c
>> +++ gcc/c/c-parser.c
>> @@ -11607,6 +11607,8 @@ c_parser_oacc_clause_async (c_parser *parser,
>> tree list)
>> static tree
>> c_parser_oacc_clause_bind (c_parser *parser, tree list)
>> {
>> + check_no_duplicate_clause (list, OMP_CLAUSE_BIND, "bind");
>> +
>> location_t loc = c_parser_peek_token (parser)->location;
>>
>> parser->lex_untranslated_string = true;
>> @@ -11615,20 +11617,43 @@ c_parser_oacc_clause_bind (c_parser *parser,
>> tree list)
>> parser->lex_untranslated_string = false;
>> return list;
>> }
>> - if (c_parser_next_token_is (parser, CPP_NAME)
>> - || c_parser_next_token_is (parser, CPP_STRING))
>> + tree name = error_mark_node;
>> + c_token *token = c_parser_peek_token (parser);
>> + if (c_parser_next_token_is (parser, CPP_NAME))
>> {
>> - tree t = c_parser_peek_token (parser)->value;
>> + tree decl = lookup_name (token->value);
>> + if (!decl)
>> + error_at (token->location, "%qE has not been declared",
>> + token->value);
>> + else if (TREE_CODE (decl) != FUNCTION_DECL)
>> + error_at (token->location, "%qE does not refer to a function",
>> + token->value);
>>
>> Quite possibly we'll want to add more error checking (matching signature
>> of X and Y, for example).
>
> Good idea, but I wonder if that would be too strict. Should we allow
> integer promotion in the bind function arguments?
I decided to be strict here. In c++ there is a decls_match function
which determines if, say, a function prototype matches it's definition.
Turns out that function was a little too strict and generic for the bind
clause, because I think we want to allow the user to bind functions
declared in different namespaces. E.g.
namespace foo {
#pragma acc routine
int bar ();
}
#pragma acc routine bind (foo::bar)
...
This should be acceptable. As a consequence, I created a
bind_decls_match function, for this purpose. I probably could have
taught the existing decls_match function how to copy with bind clauses,
but I suspect we may need to add more special cases for the bind clause
because the spec is so vague.
>> Again simplifying the c_head/clauses handling (snipped), the C++ front
>> end changes are very similar to the C front end changes:
>>
>> --- gcc/cp/parser.c
>> +++ gcc/cp/parser.c
>> @@ -31539,42 +31538,76 @@ static tree
>> cp_parser_oacc_clause_bind (cp_parser *parser, tree list)
>> {
>> [...]
>> - if (cp_lexer_next_token_is (parser->lexer, CPP_NAME)
>> - || cp_lexer_next_token_is (parser->lexer, CPP_STRING))
>> + tree name = error_mark_node;
>> + cp_token *token = cp_lexer_peek_token (parser->lexer);
>> + if (cp_lexer_next_token_is (parser->lexer, CPP_NAME))
>>
>> I'm not particularly confident in the following lookup/error checking
>> (which I copied a lot from C++ OpenACC routine parsing):
>>
>> {
>> - tree t;
>> -
>> - if (cp_lexer_peek_token (parser->lexer)->type == CPP_STRING)
>> - {
>> - t = cp_lexer_peek_token (parser->lexer)->u.value;
>> - cp_lexer_consume_token (parser->lexer);
>> + //TODO
>> + tree id = cp_parser_id_expression (parser, /*template_p=*/false,
>> + /*check_dependency_p=*/true,
>> + /*template_p=*/NULL,
>> + /*declarator_p=*/false,
>> + /*optional_p=*/false);
>> + tree decl = cp_parser_lookup_name_simple (parser, id,
>> token->location);
>> + 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)
>> + error_at (token->location, "%qE has not been declared",
>> + token->u.value);
>> + else if (/* TODO */ 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);
>> }
This namespace check is happening too early, and I don't think it's
necessary either.
>> + else if (TREE_CODE (decl) != FUNCTION_DECL)
>> + error_at (token->location,
>> + "%qE does not refer to a function",
>> + token->u.value);
>>
>> ... also we'll want to add a lot more testsuite coverage for this. (Also
>> for the OpenACC routine directive itself.)
I added some more test coverage, and I found and fixed a minor bug in
cp_finalize_oacc_routine.
>> else
>> - t = cp_parser_id_expression (parser, /*template_p=*/false,
>> - /*check_dependency_p=*/true,
>> - /*template_p=*/NULL,
>> - /*declarator_p=*/false,
>> - /*optional_p=*/false);
>> - if (t == error_mark_node)
>> - return t;
>> -
>> - tree c = build_omp_clause (loc, OMP_CLAUSE_BIND);
>> - OMP_CLAUSE_BIND_NAME (c) = t;
>> - OMP_CLAUSE_CHAIN (c) = list;
>> - list = c;
>> + {
>> + //TODO? TREE_USED (decl) = 1;
>> + tree name_id = DECL_NAME (decl);
>> + name = build_string (IDENTIFIER_LENGTH (name_id),
>> + IDENTIFIER_POINTER (name_id));
>>
>> We probably need to apply C++ name mangling here? How to do that?
Yes. But I deferred this to cp_finalize_oacc_routine because
cp_parser_oacc_clause_bind doesn't have access to the function decl.
Later on, the function decl gets mangled by calling decl_assembler_name.
>> + }
>> + //cp_lexer_consume_token (parser->lexer);
>> + }
>> + else if (cp_lexer_next_token_is (parser->lexer, CPP_STRING))
>> + {
>> + name = token->u.value;
>> + cp_lexer_consume_token (parser->lexer);
I added a new check over here for empty strings, e.g. bind (""). Are
there any other obviously invalid strings?
>> }
>> else
>> - cp_parser_error (parser, "expected identifier or character string
>> literal");
>> + cp_parser_error (parser,
>> + "expected identifier or character string literal");
>> parser->translate_strings_p = save_translate_strings_p;
>> cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN);
>> + if (name != error_mark_node)
>> + {
>> + tree c = build_omp_clause (loc, OMP_CLAUSE_BIND);
>> + OMP_CLAUSE_BIND_NAME (c) = name;
>> + OMP_CLAUSE_CHAIN (c) = list;
>> + list = c;
>> + }
>> return list;
>> }
>> --- gcc/omp-low.c
>> +++ gcc/omp-low.c
>> @@ -19853,14 +19857,76 @@ default_goacc_reduction (gcall *call)
>> static unsigned int
>> execute_oacc_device_lower ()
>> {
>> - tree attrs = get_oacc_fn_attrib (current_function_decl);
>> - int dims[GOMP_DIM_MAX];
>> -
>> - if (!attrs)
>> + /* There are offloaded functions without an "omp declare target"
>> attribute,
>> + so we'll not handle these here, but on the other hand, OpenACC
>> bind and
>> + nohost clauses can only be generated in the front ends, and an "omp
>> + declare target" attribute will then also always have been set
>> there, so
>> + this is not a problem in practice. */
>> + tree attr = lookup_attribute ("omp declare target",
>> + DECL_ATTRIBUTES (current_function_decl));
>> +
>> +#if defined(ACCEL_COMPILER)
>> + /* In an offload compiler, discard any offloaded function X that is
>> tagged
>> + with an OpenACC bind(Y) clause: all references to X have been
>> rewritten to
>> + refer to Y; X is unreachable, do not compile it. */
>> + if (attr)
>> + {
>> + tree clauses = TREE_VALUE (attr);
>> + /* TODO: device_type handling. */
>> + tree clause_bind = find_omp_clause (clauses, OMP_CLAUSE_BIND);
>> + if (clause_bind)
>> + {
>> + tree clause_bind_name = OMP_CLAUSE_BIND_NAME (clause_bind);
>> + const char *bind_name = TREE_STRING_POINTER(clause_bind_name);
>> + if (dump_file)
>> + fprintf (dump_file,
>> + "Discarding function \"%s\" with \"bind(%s)\"
>> clause.\n",
>> + IDENTIFIER_POINTER (DECL_NAME
>> (current_function_decl)),
>> + bind_name);
>> + TREE_ASM_WRITTEN (current_function_decl) = 1;
>> + return TODO_discard_function;
>> + }
>> + }
>> +#endif /* ACCEL_COMPILER */
>> +#if !defined(ACCEL_COMPILER)
>> + /* In the host compiler, discard any offloaded function that is
>> tagged with
>> + an OpenACC nohost clause. */
>> + if (attr)
>> + {
>> + tree clauses = TREE_VALUE (attr);
>> + if (find_omp_clause (clauses, OMP_CLAUSE_NOHOST))
>> + {
>> + /* There are no construct/clause combinations that could make
>> this
>> + happen, but play it safe, and verify that we never discard a
>> + function that is stored in offload_funcs, used for
>> target/offload
>> + function mapping. */
>> + if (flag_checking)
>> + {
>> + bool found = false;
>> + for (unsigned i = 0;
>> + !found && i < vec_safe_length (offload_funcs);
>> + i++)
>> + if ((*offload_funcs)[i] == current_function_decl)
>> + found = true;
>> + gcc_assert (!found);
>> + }
>> +
>> + if (dump_file)
>> + fprintf (dump_file,
>> + "Discarding function \"%s\" with \"nohost\"
>> clause.\n",
>> + IDENTIFIER_POINTER (DECL_NAME
>> (current_function_decl)));
>> + TREE_ASM_WRITTEN (current_function_decl) = 1;
>> + return TODO_discard_function;
>
> I don't think this is a good idea. If you have a nohost function,
> wounldn't that prevent the code from linking?
>
> Perhaps nohost should kind of implement a reverse bind on the host. E.g.
> discard the function defintion and replace it with an asm alias to some
> libgomp function like goacc_nohost_fallback. That way, the program will
> still link and the runtime will provide the end user with a sensible
> error when things go wrong.
Are you planning on working on this?
Note that my test coverage is limited to compile-only tests for now. I'd
like to test nohost and bind at runtime, but at least the nohosts tests
depend on this change. Do you want to work on the runtime tests?
I'll apply this patch to gomp4 in a couple of days. In the meantime,
glance at the error messages and see if they are sufficiently clear. Or
maybe I'm missing some test coverage.
Cesar
2015-12-16 Cesar Philippidis <ce...@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.
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;
+ }
+};
