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.

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;
+  }
+};
