diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 59dcc42..2d61cea 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,17 @@
+2015-12-08  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* gimplify.c (gimplify_scan_omp_clauses)
+	(gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_BIND,
+	OMP_CLAUSE_NOHOST.
+	* tree-nested.c (convert_nonlocal_omp_clauses)
+	(convert_local_omp_clauses): Likewise.
+	* omp-low.c (scan_sharing_clauses): Likewise.
+	(execute_oacc_device_lower) [ACCEL_COMPILER]: Handle OpenACC bind
+	clauses.
+	[!ACCEL_COMPILER]: Handle OpenACC nohost clauses.
+	* tree-core.h (enum omp_clause_code) <OMP_CLAUSE_BIND>: Update
+	description.
+
 2015-12-05  Chung-Lin Tang  <cltang@codesourcery.com>
 
 	* c-family/c-omp.c (c_finish_oacc_wait): Remove add_stmt() call.
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index 4701ae7..7f4e4a7 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2015-12-08  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-parser.c (c_parser_oacc_clause_bind, c_parser_oacc_routine)
+	(c_finish_oacc_routine): Update.
+
 2015-11-12  Nathan Sidwell  <nathan@codesourcery.com>
 
 	* c-typeck.c (c_finish_omp_clauses): Adjust omp_mappable_type calls.
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 14e21f5..44be0fa 100644
--- 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);
+      else
+	{
+	  //TODO? TREE_USED (decl) = 1;
+	  tree name_id = DECL_NAME (decl);
+	  name = build_string (IDENTIFIER_LENGTH (name_id),
+			       IDENTIFIER_POINTER (name_id));
+	}
       c_parser_consume_token (parser);
+    }
+  else if (c_parser_next_token_is (parser, CPP_STRING))
+    {
+      name = token->value;
+      c_parser_consume_token (parser);
+    }
+  else
+    c_parser_error (parser,
+		    "expected identifier or character string literal");
+  parser->lex_untranslated_string = false;
+  c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>");
+  if (name != error_mark_node)
+    {
       tree c = build_omp_clause (loc, OMP_CLAUSE_BIND);
-      OMP_CLAUSE_BIND_NAME (c) = t;
+      OMP_CLAUSE_BIND_NAME (c) = name;
       OMP_CLAUSE_CHAIN (c) = list;
       list = c;
     }
-  else
-    c_parser_error (parser, "expected identifier or character string literal");
-  parser->lex_untranslated_string = false;
-  c_parser_require (parser, CPP_CLOSE_PAREN, "expected %<)%>");
   return list;
 }
 
@@ -13977,10 +14002,10 @@ static void
 c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
 {
   tree decl = NULL_TREE;
-  /* Create a dummy claue, to record location.  */
+  /* Create a dummy clause, to record the location.  */
   tree c_head = build_omp_clause (c_parser_peek_token (parser)->location,
-				  OMP_CLAUSE_SEQ);
-  
+				  OMP_CLAUSE_ERROR);
+
   if (context != pragma_external)
     c_parser_error (parser, "%<#pragma acc routine%> not at file scope");
 
@@ -14018,9 +14043,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
   tree clauses = c_parser_oacc_all_clauses
     (parser, OACC_ROUTINE_CLAUSE_MASK, "#pragma acc routine",
      OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK);
-
-  /* Force clauses to be non-null, by attaching context to it.  */
-  clauses = tree_cons (c_head, clauses, NULL_TREE);
+  /* Prepend the dummy clause.  */
+  OMP_CLAUSE_CHAIN (c_head) = clauses;
+  clauses = c_head;
   
   if (decl)
     c_finish_oacc_routine (parser, decl, clauses, true, true, false);
@@ -14040,7 +14065,9 @@ static void
 c_finish_oacc_routine (c_parser *ARG_UNUSED (parser), tree fndecl,
 		       tree clauses, bool named, bool first, bool is_defn)
 {
-  location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses));
+  location_t loc = OMP_CLAUSE_LOCATION (clauses);
+  /* Get rid of the dummy clause.  */
+  clauses = OMP_CLAUSE_CHAIN (clauses);
 
   if (!fndecl || TREE_CODE (fndecl) != FUNCTION_DECL || !first)
     {
@@ -14059,13 +14086,12 @@ c_finish_oacc_routine (c_parser *ARG_UNUSED (parser), tree fndecl,
 	      TREE_USED (fndecl) ? "use" : "definition");
 
   /* Process for function attrib  */
-  tree dims = build_oacc_routine_dims (TREE_VALUE (clauses));
+  tree dims = build_oacc_routine_dims (clauses);
   replace_oacc_fn_attrib (fndecl, dims);
 
-  /* Also attach as a declare.  */
-  DECL_ATTRIBUTES (fndecl)
-    = tree_cons (get_identifier ("omp declare target"),
-		 clauses, DECL_ATTRIBUTES (fndecl));
+  /* Also add an "omp declare target" attribute, with clauses.  */
+  DECL_ATTRIBUTES (fndecl) = tree_cons (get_identifier ("omp declare target"),
+					clauses, DECL_ATTRIBUTES (fndecl));
 }
 
 /* OpenACC 2.0:
diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp
index e4d000d..3f1f37e 100644
--- gcc/cp/ChangeLog.gomp
+++ gcc/cp/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2015-12-08  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* parser.c (cp_ensure_no_oacc_routine, cp_parser_oacc_clause_bind)
+	(cp_parser_oacc_routine, cp_parser_late_parsing_oacc_routine)
+	(cp_finalize_oacc_routine): Update.
+	* pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_BIND,
+	OMP_CLAUSE_NOHOST.
+
 2015-11-12  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* semantics.c (finish_omp_clauses): Remove "reference types are
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 9d18cfc..6556db3 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -1326,10 +1326,9 @@ cp_ensure_no_oacc_routine (cp_parser *parser)
 {
   if (parser->oacc_routine && !parser->oacc_routine->error_seen)
     {
-      tree clauses = parser->oacc_routine->clauses;
-      location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE (clauses));
-
-      error_at (loc, "%<#pragma oacc routine%> not followed by function "
+      /* The first clause is a dummy, providing location information.  */
+      error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses),
+		"%<#pragma oacc routine%> not followed by function "
 		"declaration or definition");
       parser->oacc_routine = NULL;
     }
@@ -31539,42 +31538,76 @@ static tree
 cp_parser_oacc_clause_bind (cp_parser *parser, tree list)
 {
   location_t loc = cp_lexer_peek_token (parser->lexer)->location;
+
+  check_no_duplicate_clause (list, OMP_CLAUSE_BIND, "bind", loc);
+
   bool save_translate_strings_p = parser->translate_strings_p;
-
   parser->translate_strings_p = false;
   if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
     {
       parser->translate_strings_p = save_translate_strings_p;
       return 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))
     {
-      tree t;
-
-      if (cp_lexer_peek_token (parser->lexer)->type == CPP_STRING)
+      //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))
 	{
-	  t = cp_lexer_peek_token (parser->lexer)->u.value;
-	  cp_lexer_consume_token (parser->lexer);
+	  /* 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
-	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;
-
+	{
+	  //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);
+    }
+  else if (cp_lexer_next_token_is (parser->lexer, CPP_STRING))
+    {
+      name = token->u.value;
+      cp_lexer_consume_token (parser->lexer);
+    }
+  else
+    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) = t;
+      OMP_CLAUSE_BIND_NAME (c) = name;
       OMP_CLAUSE_CHAIN (c) = list;
       list = c;
     }
-  else
-    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);
   return list;
 }
 
@@ -36020,9 +36053,8 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
       parser->oacc_routine = &data;
     }
 
-  tree decl = NULL_TREE;
-  /* Create a dummy claue, to record location.  */
-  tree c_head = build_omp_clause (pragma_tok->location, OMP_CLAUSE_SEQ);
+  /* Create a dummy clause, to record the location.  */
+  tree c_head = build_omp_clause (pragma_tok->location, OMP_CLAUSE_ERROR);
 
   if (context != pragma_external)
     {
@@ -36044,6 +36076,7 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
 	    parser->oacc_routine->error_seen = true;
 	  cp_parser_require_pragma_eol (parser, pragma_tok);
 
+	  /* The first clause is a dummy, providing location information.  */
 	  error_at (OMP_CLAUSE_LOCATION (parser->oacc_routine->clauses),
 		    "%<#pragma oacc routine%> not followed by a single "
 		    "function declaration or definition");
@@ -36064,7 +36097,7 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
 					 /*template_p=*/NULL,
 					 /*declarator_p=*/false,
 					 /*optional_p=*/false);
-      decl = cp_parser_lookup_name_simple (parser, id, token->location);
+      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);
@@ -36079,14 +36112,14 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
 
       /* Build a chain of clauses.  */
       parser->lexer->in_pragma = true;
-      tree clauses = NULL_TREE;
-      clauses = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
-					    "#pragma acc routine",
-					    cp_lexer_peek_token
-					    (parser->lexer));
+      tree clauses
+	= cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
+				      "#pragma acc routine",
+				      cp_lexer_peek_token (parser->lexer));
 
-      /* Force clauses to be non-null, by attaching context to it.  */
-      clauses = tree_cons (c_head, clauses, NULL_TREE);
+      /* Prepend the dummy clause.  */
+      OMP_CLAUSE_CHAIN (c_head) = clauses;
+      clauses = c_head;
 
       if (decl && is_overloaded_fn (decl)
 	  && (TREE_CODE (decl) != FUNCTION_DECL
@@ -36142,9 +36175,7 @@ cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
 
       if (first_p)
 	{
-	  /* Create an empty list of clauses.  */
-	  parser->oacc_routine->clauses = tree_cons (c_head, NULL_TREE,
-						     NULL_TREE);
+	  parser->oacc_routine->clauses = c_head;
 	  cp_parser_declaration (parser);
 
 	  if (parser->oacc_routine
@@ -36168,10 +36199,12 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
   struct cp_token_cache *ce;
   cp_omp_declare_simd_data *data = parser->oacc_routine;
   tree cl, clauses = parser->oacc_routine->clauses;
-  location_t loc;
 
-  loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses));
-  
+  /* The first clause is a dummy, providing location information.  */
+  location_t loc = OMP_CLAUSE_LOCATION (clauses);
+  /* Get rid of it now.  */
+  clauses = OMP_CLAUSE_CHAIN (clauses);
+
   if ((!data->error_seen && data->fndecl_seen)
       || data->tokens.length () != 1)
     {
@@ -36195,10 +36228,12 @@ cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
 				   OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK);
   cp_parser_pop_lexer (parser);
 
-  tree c_head = build_omp_clause (loc, OMP_CLAUSE_SEQ);
+  /* Create a dummy clause, to record the location.  */
+  tree c_head = build_omp_clause (loc, OMP_CLAUSE_ERROR);
 
-  /* Force clauses to be non-null, by attaching context to it.  */
-  parser->oacc_routine->clauses = tree_cons (c_head, cl, NULL_TREE);
+  /* Prepend the dummy clause.  */
+  OMP_CLAUSE_CHAIN (c_head) = cl;
+  parser->oacc_routine->clauses = c_head;
 
   data->fndecl_seen = true;
   return attrs;
@@ -36213,7 +36248,9 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
   if (__builtin_expect (parser->oacc_routine != NULL, 0))
     {
       tree clauses = parser->oacc_routine->clauses;
-      location_t loc = OMP_CLAUSE_LOCATION (TREE_PURPOSE(clauses));
+      location_t loc = OMP_CLAUSE_LOCATION (clauses);
+      /* Get rid of the dummy clause.  */
+      clauses = OMP_CLAUSE_CHAIN (clauses);
 
       if (parser->oacc_routine->error_seen)
 	return;
@@ -36252,13 +36289,13 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
 	}
 
       /* Process for function attrib  */
-      tree dims = build_oacc_routine_dims (TREE_VALUE (clauses));
+      tree dims = build_oacc_routine_dims (clauses);
       replace_oacc_fn_attrib (fndecl, dims);
       
-      /* Add an "omp target" attribute.  */
+      /* Also add an "omp declare target" attribute, with clauses.  */
       DECL_ATTRIBUTES (fndecl)
 	= tree_cons (get_identifier ("omp declare target"),
-		     NULL_TREE, DECL_ATTRIBUTES (fndecl));
+		     clauses, DECL_ATTRIBUTES (fndecl));
     }
 }
 
diff --git gcc/cp/pt.c gcc/cp/pt.c
index 93f6e6d..0d2fe63 100644
--- gcc/cp/pt.c
+++ gcc/cp/pt.c
@@ -14501,6 +14501,8 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
 	      }
 	  }
 	  break;
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp
index 6c7b8af..00e5746 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,3 +1,12 @@
+2015-12-08  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* gfortran.h (symbol_attribute): Add oacc_function_nohost member.
+	* openmp.c (gfc_match_oacc_routine): Set it.
+	* trans-decl.c (add_attributes_to_decl): Use it to decide whether
+	to generate an OMP_CLAUSE_NOHOST clause.
+	* trans-openmp.c (gfc_trans_omp_clauses_1): Unreachable code to
+	generate an OMP_CLAUSE_NOHOST clause.
+
 2015-12-03  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* openmp.c (gfc_match_omp_clauses): Allow subarrays for acc reductions.
diff --git gcc/fortran/gfortran.h gcc/fortran/gfortran.h
index 26f4c8a..2c8c806 100644
--- gcc/fortran/gfortran.h
+++ gcc/fortran/gfortran.h
@@ -850,6 +850,7 @@ typedef struct
 
   /* This is an OpenACC acclerator function at level N - 1  */
   unsigned oacc_function:3;
+  unsigned oacc_function_nohost:1;
 
   /* Attributes set by compiler extensions (!GCC$ ATTRIBUTES).  */
   unsigned ext_attr:EXT_ATTR_NUM;
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index e7f61f2..b59528be 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -1884,6 +1884,8 @@ gfc_match_oacc_routine (void)
 	goto cleanup;
       gfc_current_ns->proc_name->attr.oacc_function
 	= gfc_oacc_routine_dims (c) + 1;
+      gfc_current_ns->proc_name->attr.oacc_function_nohost
+	= c ? c->nohost : false;
     }
 
   if (n)
diff --git gcc/fortran/trans-decl.c gcc/fortran/trans-decl.c
index eaf46cb..2fe4abd 100644
--- gcc/fortran/trans-decl.c
+++ gcc/fortran/trans-decl.c
@@ -1309,8 +1309,13 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list)
       || sym_attr.oacc_declare_device_resident
 #endif
       )
-    list = tree_cons (get_identifier ("omp declare target"),
-		      NULL_TREE, list);
+    {
+      tree c = NULL_TREE;
+      if (sym_attr.oacc_function_nohost)
+	c = build_omp_clause (/* TODO */ input_location,
+			      OMP_CLAUSE_NOHOST);
+      list = tree_cons (get_identifier ("omp declare target"), c, list);
+    }
 #if 0 /* TODO */
   if (sym_attr.oacc_declare_link)
     list = tree_cons (get_identifier ("omp declare target link"),
diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c
index 6ed4a57..4de4726 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -2644,6 +2644,13 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
 	  OMP_CLAUSE_GANG_STATIC_EXPR (c) = arg;
 	}
     }
+  if (clauses->nohost)
+    {
+      c = build_omp_clause (where.lb->location, OMP_CLAUSE_NOHOST);
+      omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+      //TODO
+      gcc_unreachable();
+    }
 
   return nreverse (omp_clauses);
 }
diff --git gcc/gimplify.c gcc/gimplify.c
index b00de81..e8964c6 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -7413,6 +7413,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
 	  break;
 
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
@@ -8104,6 +8106,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	case OMP_CLAUSE_DEVICE_TYPE:
 	  break;
 
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
diff --git gcc/lto/ChangeLog.gomp gcc/lto/ChangeLog.gomp
index 03ed7b7..635bdfa 100644
--- gcc/lto/ChangeLog.gomp
+++ gcc/lto/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2015-12-08  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* lto.c (read_cgraph_and_symbols) [ACCEL_COMPILER]: Handle OpenACC
+	bind clauses.
+
 2015-08-31  Nathan Sidwell  <nathan@codesourcery.com>
 
 	* lto-lang.c (DEF_FUNCTION_TYPE_VAR_6): Define.
diff --git gcc/lto/lto.c gcc/lto/lto.c
index b1e2d6e..5820feb 100644
--- gcc/lto/lto.c
+++ gcc/lto/lto.c
@@ -49,6 +49,8 @@ along with GCC; see the file COPYING3.  If not see
 #include "params.h"
 #include "ipa-utils.h"
 #include "gomp-constants.h"
+#include "omp-low.h"
+#include "stringpool.h"
 
 
 /* Number of parallel tasks to run, -1 if we want to use GNU Make jobserver.  */
@@ -2942,6 +2944,36 @@ read_cgraph_and_symbols (unsigned nfiles, const char **fnames)
 
   ggc_free (all_file_decl_data);
   all_file_decl_data = NULL;
+
+#ifdef ACCEL_COMPILER
+  /* In an offload compiler, redirect calls to any function X that is tagged
+     with an OpenACC bind(Y) clause to call Y instead of X.  */
+  FOR_EACH_SYMBOL (snode)
+  {
+    tree decl = snode->decl;
+    tree attr = lookup_attribute ("omp declare target",
+				  DECL_ATTRIBUTES (decl));
+    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 (symtab->dump_file)
+	      fprintf (symtab->dump_file,
+		       "Applying \"bind(%s)\" clause to declaration of "
+		       "function \"%s\".\n",
+		       bind_name, IDENTIFIER_POINTER (DECL_NAME (decl)));
+	    //TODO: Use gcc/varasm.c:set_user_assembler_name instead?
+	    symtab->change_decl_assembler_name (decl,
+						get_identifier (bind_name));
+	  }
+      }
+  }
+#endif /* ACCEL_COMPILER */
 }
 
 
diff --git gcc/omp-low.c gcc/omp-low.c
index 88e41b8..9ef7161 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -2279,6 +2279,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  sorry ("Clause not supported yet");
 	  break;
 
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
@@ -2453,6 +2455,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  sorry ("Clause not supported yet");
 	  break;
 
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
@@ -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;
+	}
+    }
+#endif /* !ACCEL_COMPILER */
+
+  attr = get_oacc_fn_attrib (current_function_decl);
+  if (!attr)
     /* Not an offloaded function.  */
     return 0;
-
-  int fn_level = oacc_validate_dims (current_function_decl, attrs, dims);
+  int dims[GOMP_DIM_MAX];
+  int fn_level = oacc_validate_dims (current_function_decl, attr, dims);
 
   /* Discover, partition and process the loops.  */
   oacc_loop *loops = oacc_loop_discovery ();
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 1135ce0..de3a68a 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,14 @@
+2015-12-08  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/routine-1.c: Update.
+	* c-c++-common/goacc/routine-2.c: Likewise.
+	* c-c++-common/goacc/routine-5.c: Likewise.
+	* c-c++-common/goacc/routine-8.c: Remove file.
+	* c-c++-common/goacc/routine-9.c: Remove file.
+	* c-c++-common/goacc/routine-nohost-1.c: New file.
+	* g++.dg/goacc/routine-1.C: Likewise.
+	* g++.dg/goacc/routine-2.C: Likewise.
+
 2015-12-03  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* gfortran.dg/goacc/array-reduction.f90: New test.
diff --git gcc/testsuite/c-c++-common/goacc/routine-1.c gcc/testsuite/c-c++-common/goacc/routine-1.c
index a5e0d69..6535c8c 100644
--- gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -1,3 +1,4 @@
+/* Test valid use of clauses with routine.  */
 
 #pragma acc routine gang
 void gang (void)
@@ -19,15 +20,45 @@ void seq (void)
 {
 }
 
+#pragma acc routine
+void bind_f_1 (void)
+{
+}
+
+#pragma acc routine bind (bind_f_1)
+void bind_f_1_1 (void)
+{
+}
+
+/* Non-sensical bind clause, but permitted.  */
+#pragma acc routine bind ("bind_f_2")
+void bind_f_2 (void)
+{
+}
+
+#pragma acc routine bind ("bind_f_2")
+void bind_f_2_1 (void)
+{
+}
+
+#pragma acc routine nohost
+void nohost (void)
+{
+}
+
 int main ()
 {
-
 #pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
   {
     gang ();
     worker ();
     vector ();
     seq ();
+    bind_f_1 ();
+    bind_f_1_1 ();
+    bind_f_2 ();
+    bind_f_2_1 ();
+    nohost ();
   }
 
   return 0;
diff --git gcc/testsuite/c-c++-common/goacc/routine-2.c gcc/testsuite/c-c++-common/goacc/routine-2.c
index fc5eb11..35857ea 100644
--- gcc/testsuite/c-c++-common/goacc/routine-2.c
+++ gcc/testsuite/c-c++-common/goacc/routine-2.c
@@ -1,21 +1,40 @@
+/* Test invalid use of clauses with routine.  */
+
 #pragma acc routine gang worker /* { dg-error "multiple loop axes" } */
-void gang (void)
-{
-}
+extern void gang (void);
 
 #pragma acc routine worker vector /* { dg-error "multiple loop axes" } */
-void worker (void)
-{
-}
+extern void worker (void);
 
 #pragma acc routine vector seq /* { dg-error "multiple loop axes" } */
-void vector (void)
-{
-}
+extern void vector (void);
 
 #pragma acc routine seq gang /* { dg-error "multiple loop axes" } */
-void seq (void)
-{
-}
+extern void seq (void);
 
-#pragma acc routine (nothing) gang /* { dg-error "not been declared" } */
+extern float F;
+#pragma acc routine bind (F) /* { dg-error ".F. does not refer to a function" } */
+extern void F_1 (void);
+
+typedef int T;
+#pragma acc routine bind (T) /* { dg-error ".T. does not refer to a function" } */
+extern void T_1 (void);
+
+#pragma acc routine (nothing) gang /* { dg-error ".nothing. has not been declared" } */
+
+#pragma acc routine bind (bind_0) /* { dg-error ".bind_0. has not been declared" }*/
+extern void bind_0 (void);
+
+extern void a(void), b(void);
+
+#pragma acc routine bind(a) bind(b) /* { dg-error "too many .bind. clauses" } */
+extern void bind_1 (void);
+
+#pragma acc routine bind(a) bind("b") /* { dg-error "too many .bind. clauses" } */
+extern void bind_2 (void);
+
+#pragma acc routine bind("a") bind("b") /* { dg-error "too many .bind. clauses" } */
+extern void bind_3 (void);
+
+#pragma acc routine nohost nohost /* { dg-error "too many .nohost. clauses" } */
+extern void nohost (void);
diff --git gcc/testsuite/c-c++-common/goacc/routine-5.c gcc/testsuite/c-c++-common/goacc/routine-5.c
index ccda097..f4ae843 100644
--- gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -45,3 +45,17 @@ using namespace g;
 #pragma acc routine (a) /* { dg-error "does not refer to" } */
   
 #pragma acc routine (c) /* { dg-error "does not refer to" } */
+
+
+void Bar ();
+
+void Foo ()
+{
+  Bar ();
+}
+
+#pragma acc routine (Bar) // { dg-error "must be applied before use" }
+
+#pragma acc routine (Foo) gang // { dg-error "must be applied before definition" }
+
+#pragma acc routine (Baz) // { dg-error "not been declared" }
diff --git gcc/testsuite/c-c++-common/goacc/routine-8.c gcc/testsuite/c-c++-common/goacc/routine-8.c
deleted file mode 100644
index e35dfc1..0000000
--- gcc/testsuite/c-c++-common/goacc/routine-8.c
+++ /dev/null
@@ -1,52 +0,0 @@
-/* Test valid use of clauses with routine.  */
-/* { dg-do compile } */
-
-#pragma acc routine gang
-void
-f1 (void)
-{
-}
-
-#pragma acc routine worker
-void
-f2 (void)
-{
-}
-
-#pragma acc routine vector
-void
-f3 (void)
-{
-}
-
-#pragma acc routine seq
-void
-f4 (void)
-{
-}
-
-#pragma acc routine bind (f4a)
-void
-f5 (void)
-{
-}
-
-typedef int T;
-
-#pragma acc routine bind (T)
-void
-f6 (void)
-{
-}
-
-#pragma acc routine bind ("f7a")
-void
-f7 (void)
-{
-}
-
-#pragma acc routine nohost
-void
-f8 (void)
-{
-}
diff --git gcc/testsuite/c-c++-common/goacc/routine-9.c gcc/testsuite/c-c++-common/goacc/routine-9.c
deleted file mode 100644
index f712a6f..0000000
--- gcc/testsuite/c-c++-common/goacc/routine-9.c
+++ /dev/null
@@ -1,20 +0,0 @@
-/* Test invalid use of clauses with routine.  */
-/* { dg-do compile } */
-
-void Bar ();
-
-void Foo ()
-{
-  Bar ();
-}
-
-#pragma acc routine (Bar) // { dg-error "must be applied before use" }
-
-#pragma acc routine (Foo) gang // { dg-error "must be applied before definition" }
-
-#pragma acc routine (Baz) // { dg-error "not been declared" }
-
-#pragma acc routine
-int i;
-// { dg-error "not followed by single function" "" { target c } 17 }
-// { dg-error "not followed by function declaration or definition" "" { target c++ } 17 }
diff --git gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
new file mode 100644
index 0000000..88af656
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
@@ -0,0 +1,34 @@
+/* Test the nohost clause for OpenACC routine directive.  Exercising different
+   variants for declaring routines.  */
+
+/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+#pragma acc routine nohost
+int THREE(void)
+{
+  return 3;
+}
+
+/* { dg-final { scan-tree-dump "Discarding function .THREE. with .nohost. clause" "oaccdevlow" } } */
+
+
+#pragma acc routine nohost
+extern void NOTHING(void);
+
+void NOTHING(void)
+{
+}
+
+/* { dg-final { scan-tree-dump "Discarding function .NOTHING. with .nohost. clause" "oaccdevlow" } } */
+
+
+extern float ADD(float, float);
+
+#pragma acc routine (ADD) nohost
+
+float ADD(float x, float y)
+{
+  return x + y;
+}
+
+/* { dg-final { scan-tree-dump "Discarding function .ADD. with .nohost. clause" "oaccdevlow" } } */
diff --git gcc/testsuite/g++.dg/goacc/routine-1.C gcc/testsuite/g++.dg/goacc/routine-1.C
new file mode 100644
index 0000000..a73a73d
--- /dev/null
+++ gcc/testsuite/g++.dg/goacc/routine-1.C
@@ -0,0 +1,13 @@
+/* Test valid use of the routine directive.  */
+
+namespace N
+{
+  extern void foo1();
+  extern void foo2();
+#pragma acc routine (foo1)
+#pragma acc routine
+  void foo3()
+  {
+  }
+}
+#pragma acc routine (N::foo2)
diff --git gcc/testsuite/g++.dg/goacc/routine-2.C gcc/testsuite/g++.dg/goacc/routine-2.C
new file mode 100644
index 0000000..92fc161
--- /dev/null
+++ gcc/testsuite/g++.dg/goacc/routine-2.C
@@ -0,0 +1,16 @@
+/* Test invalid use of the routine directive.  */
+
+// { dg-do compile }
+// { dg-options "-fopenacc" }
+
+template <typename T>
+extern T one_d();
+#pragma acc routine (one_d) nohost /* { dg-error "names a set of overloads" } */
+
+template <typename T>
+T
+one()
+{
+  return 1;
+}
+#pragma acc routine (one) bind(one_d) /* { dg-error "names a set of overloads" } */
diff --git gcc/tree-core.h gcc/tree-core.h
index 46a42da..43507de 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -461,7 +461,7 @@ enum omp_clause_code {
   /* OpenACC clause: vector_length (integer-expression).  */
   OMP_CLAUSE_VECTOR_LENGTH,
 
-  /* OpenACC clause: bind ( identifer | string ).  */
+  /* OpenACC clause: bind (string).  */
   OMP_CLAUSE_BIND,
 
   /* OpenACC clause: nohost.  */
diff --git gcc/tree-nested.c gcc/tree-nested.c
index da19e8d..7198f1e 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1200,6 +1200,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_SEQ:
 	  break;
 
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
@@ -1882,6 +1884,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
 	case OMP_CLAUSE_SEQ:
 	  break;
 
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 791aa4c..a59cc9d 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2015-12-08  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c: New
+	file.
+	* testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c:
+	Likewise.
+	* testsuite/libgomp.oacc-fortran/routine-6.f90: XFAIL.
+
 2015-12-06  James Norris  <jnorris@codesourcery.com>
 
 	* oacc-parallel.c (GOACC_parallel_keyed, GOACC_data_start):
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c
new file mode 100644
index 0000000..b991bb1
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c
@@ -0,0 +1,105 @@
+/* Test the bind and nohost clauses for OpenACC routine directive.  */
+
+/* TODO.  Function inlining and the OpenACC bind clause do not yet get on well
+   with one another.
+   { dg-additional-options "-fno-inline" } */
+
+/* TODO.  C works, but for C++ we get: "lto1: internal compiler error: in
+   ipa_propagate_frequency".
+   { dg-xfail-if "TODO" { *-*-* } } */
+
+#include <openacc.h>
+
+/* "MINUS_TWO" is the device variant for function "TWO".  Similar for "THREE",
+   and "FOUR".  Exercising different variants for declaring routines.  */
+
+#pragma acc routine nohost
+extern int MINUS_TWO(void);
+
+int MINUS_TWO(void)
+{
+  if (!acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return -2;
+}
+
+extern int TWO(void);
+#pragma acc routine (TWO) bind(MINUS_TWO)
+
+int TWO(void)
+{
+  if (acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return 2;
+}
+
+
+#pragma acc routine nohost
+int MINUS_THREE(void)
+{
+  if (!acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return -3;
+}
+
+#pragma acc routine bind(MINUS_THREE)
+extern int THREE(void);
+
+int THREE(void)
+{
+  if (acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return 3;
+}
+
+
+/* Due to using a string in the bind clause, we don't need "MINUS_FOUR" in
+   scope here.  */
+#pragma acc routine bind("MINUS_FOUR")
+int FOUR(void)
+{
+  if (acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return 4;
+}
+
+extern int MINUS_FOUR(void);
+#pragma acc routine (MINUS_FOUR) nohost
+
+int MINUS_FOUR(void)
+{
+  if (!acc_on_device(acc_device_not_host))
+    __builtin_abort();
+  return -4;
+}
+
+
+int main()
+{
+  int x2, x3, x4;
+
+#pragma acc parallel copyout(x2, x3, x4) if(0)
+  {
+    x2 = TWO();
+    x3 = THREE();
+    x4 = FOUR();
+  }
+  if (x2 != 2 || x3 != 3 || x4 != 4)
+    __builtin_abort();
+
+#pragma acc parallel copyout(x2, x3, x4)
+  {
+    x2 = TWO();
+    x3 = THREE();
+    x4 = FOUR();
+  }
+#ifdef ACC_DEVICE_TYPE_host
+  if (x2 != 2 || x3 != 3 || x4 != 4)
+    __builtin_abort();
+#else
+  if (x2 != -2 || x3 != -3 || x4 != -4)
+    __builtin_abort();
+#endif
+
+  return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
new file mode 100644
index 0000000..365af93
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c
@@ -0,0 +1,18 @@
+/* { dg-do link } */
+
+extern int three (void);
+
+#pragma acc routine (three) nohost
+__attribute__((noinline))
+int three(void)
+{
+  return 3;
+}
+
+int main(void)
+{
+  return (three() == 3) ? 0 : 1;
+}
+
+/* Expecting link to fail; "undefined reference to `three'" (or similar).
+   { dg-excess-errors "" } */
diff --git libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90 libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90
index 9ba6da8..1bae09c 100644
--- libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/routine-6.f90
@@ -1,5 +1,5 @@
 ! { dg-do run }
-! { dg-xfail-if "not found" { openacc_host_selected } }
+! { dg-xfail-if "TODO" { *-*-* } }
 
 program main
   integer :: a, n
