Hi!

On Sat, 14 Nov 2015 09:36:36 +0100, I wrote:
> Initial support for the OpenACC bind and nohost clauses (routine
> directive) for C, C++.  Fortran to follow.  Middle end handling and more
> complete testsuite coverage also to follow once we got a few details
> clarified.  OK for trunk?

(Has not yet been reviewed.)  Meanwhile, I continued working on the
implementation, focussing on C.  See also my question "How to rewrite
call targets (OpenACC bind clause)",
<http://news.gmane.org/find-root.php?message_id=%3C877fkq482i.fsf%40hertz.schwinge.homeip.net%3E>.

To enable Cesar to help with the C++ and Fortran front ends (thanks!), in
r231423, I just committed "[WIP] OpenACC bind, nohost clauses" to
gomp-4_0-branch.  (There has already been initial support, parsing only,
on gomp-4_0-branch.)  I'll try to make progress with the generic middle
end bits, but will appreciate any review comments, so before inlining the
complete patch, first a few questions/comments:

In the OpenACC bind(Y) clause attached to a routine(X) directive, Y can
be an identifier or a string.  In the front ends, I canonicalize that
into a string, as we -- at least currently -- don't have any use for the
identifier (or decl?) later on:

    --- gcc/tree-core.h
    +++ gcc/tree-core.h
    @@ -461,7 +461,7 @@ enum omp_clause_code {
    -  /* OpenACC clause: bind ( identifer | string ).  */
    +  /* OpenACC clause: bind (string).  */
       OMP_CLAUSE_BIND,

All the following are unreachable for OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST;
document that to make it obvious/expected:

    --- 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 ();
            }
    --- 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 ();
            }
    --- 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 ();
            }
    --- 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 ();
            }

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).

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

Should I set TREE_USED after having looked up the identifier?

    +  else if (c_parser_next_token_is (parser, CPP_STRING))
    +    {
    +      name = token->value;
           c_parser_consume_token (parser);
    -      tree c = build_omp_clause (loc, OMP_CLAUSE_BIND);
    -      OMP_CLAUSE_BIND_NAME (c) = t;
    -      OMP_CLAUSE_CHAIN (c) = list;
    -      list = c;
         }
       else
    -    c_parser_error (parser, "expected identifier or character string 
literal");
    +    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) = name;
    +      OMP_CLAUSE_CHAIN (c) = list;
    +      list = c;
    +    }
       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);

I don't know why somebody chose OMP_CLAUSE_SEQ for this; changed to a
distinctive OMP_CLAUSE_ERROR.  In the following, handling of c_head and
generally the clauses seemed unnecessarily complicated to me, so I
simplified that as follows:

    @@ -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));
     }

I don't know why somebody chose to attach the clauses to the "omp declare
target" attribute in this way?  Especially given that so far there hasn't
been any user of this information (I'm now adding such users).  Is that
OK, or should we have a separate "omp clauses" attribute or similar?

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

           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?

    +       }
    +      //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");
    +    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;
     }

What I changed in the Fortran front end is just a quick hack.  Also I
have not spent any effort on updating the existing OpenACC bind clause
support: the name is (only) parsed into routine_bind, but then not
handled any further?  Also needs testsuite coverage, obviously.

    --- 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;
    --- 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)
    --- 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"),

I guess add_attributes_to_decl is the correct place to be doning this?

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

Probably we can generally just put a gcc_unreachable call here, with a
source code comment added.  Again, this is to make sure that the reader
of that code doesn't wonder why "clauses->nohost" has been forgotten to
be handled here.

       return nreverse (omp_clauses);
     }

Middle end.  In the LTO wrapper, at the end of read_cgraph_and_symbols,
for ACCEL_COMPILERs handle OpenACC bind clauses:

    --- gcc/lto/lto.c
    +++ gcc/lto/lto.c
    @@ -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 */
     }

Probably that should be put into a separate function (in gcc/omp-low.c,
even?).  Is the end of read_cgraph_and_symbols the correct place to
put/call this?  Per my "How to rewrite call targets (OpenACC bind
clause)" email,
<http://news.gmane.org/find-root.php?message_id=%3C877fkq482i.fsf%40hertz.schwinge.homeip.net%3E>,
it's still not clear to me whether just setting the decl's assembler name
here is the right (and sufficient) thing to do (but it seems to work,
with -fno-inline at least...).

Joseph once pointed out that we'll need to add user_label_prefix to the
bind_name -- but only if an indentifier has been used for Y in the
bind(Y) clause, and not when a string has been used?

Then, the following handling in execute_oacc_device_lower (correct
position in the pipeline -- as early as possible after the LTO front end,
I guess?), for ACCEL_COMPILERs handle OpenACC bind clauses, and for
non-ACCEL_COMPILERs handle OpenACC nohost clauses.  In both cases, use
the new TODO_discard_function,
<http://news.gmane.org/find-root.php?message_id=%3C563A3791.7020001%40suse.cz%3E>,
that has recently been added.  :-)

    --- 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;
    +       }
    +    }
    +#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 ();

Initial testsuite updates:

    --- 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.  */
    [...]
    +extern void a(void), b(void);
    +
    +#pragma acc routine bind(a) bind(b) /* { dg-error "too many .bind. 
clauses" } */
    +extern void bind_1 (void);

This diagnostic does make sense (can't bind to a and b at the same time),
but this will need re-visiting for device_type clause support.

    +#pragma acc routine nohost nohost /* { dg-error "too many .nohost. 
clauses" } */
    +extern void nohost (void);

But I'm not too sure about this one.  After all, there is no harm in
specifying multiple such clauses.  However, GCC generally (also for
"simple" OpenMP clauses?) seems to diagnose such usage, so it's probably
a good idea to be consistent?

    --- /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.

    +/* TODO.  C works, but for C++ we get: "lto1: internal compiler error: in
    +   ipa_propagate_frequency".
    +   { dg-xfail-if "TODO" { *-*-* } } */

TODO.  Perhaps related to missing C++ name mangling (see above), perhaps
something else.

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

I'd also like to add test cases where the host and device function
definitions are in separate files, so I'll try to figure out how to do
that in the libgomp testsuite.

    --- /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 "" } */

This results in an XFAIL, which is not nice.  Is there a mechanism in the
GCC testsuite/DejaGnu to check for an expected link failure (due to a
missing symbol)?  I guess we could cook up something that instead
triggers a link failure for a duplicate or incompatible symbol
definition?

    --- 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" { *-*-* } }

TODO.  ICE, if I remember correctly.


The complete patch:

commit 00fe55db1391906f85a025faa6a5d72ad995b2b6
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Tue Dec 8 19:43:20 2015 +0000

    [WIP] OpenACC bind, nohost clauses
    
        gcc/c/
        * c-parser.c (c_parser_oacc_clause_bind, c_parser_oacc_routine)
        (c_finish_oacc_routine): Update.
        gcc/cp/
        * 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.
        gcc/fortran/
        * 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.
        gcc/
        * 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.
        gcc/lto/
        * lto.c (read_cgraph_and_symbols) [ACCEL_COMPILER]: Handle OpenACC
        bind clauses.
        gcc/testsuite/
        * 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.
        libgomp/
        * 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.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@231423 
138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp                                 |  14 +++
 gcc/c/ChangeLog.gomp                               |   5 +
 gcc/c/c-parser.c                                   |  66 +++++++---
 gcc/cp/ChangeLog.gomp                              |   8 ++
 gcc/cp/parser.c                                    | 135 +++++++++++++--------
 gcc/cp/pt.c                                        |   2 +
 gcc/fortran/ChangeLog.gomp                         |   9 ++
 gcc/fortran/gfortran.h                             |   1 +
 gcc/fortran/openmp.c                               |   2 +
 gcc/fortran/trans-decl.c                           |   9 +-
 gcc/fortran/trans-openmp.c                         |   7 ++
 gcc/gimplify.c                                     |   4 +
 gcc/lto/ChangeLog.gomp                             |   5 +
 gcc/lto/lto.c                                      |  32 +++++
 gcc/omp-low.c                                      |  78 +++++++++++-
 gcc/testsuite/ChangeLog.gomp                       |  11 ++
 gcc/testsuite/c-c++-common/goacc/routine-1.c       |  33 ++++-
 gcc/testsuite/c-c++-common/goacc/routine-2.c       |  45 +++++--
 gcc/testsuite/c-c++-common/goacc/routine-5.c       |  14 +++
 gcc/testsuite/c-c++-common/goacc/routine-8.c       |  52 --------
 gcc/testsuite/c-c++-common/goacc/routine-9.c       |  20 ---
 .../c-c++-common/goacc/routine-nohost-1.c          |  34 ++++++
 gcc/testsuite/g++.dg/goacc/routine-1.C             |  13 ++
 gcc/testsuite/g++.dg/goacc/routine-2.C             |  16 +++
 gcc/tree-core.h                                    |   2 +-
 gcc/tree-nested.c                                  |   4 +
 libgomp/ChangeLog.gomp                             |   8 ++
 .../routine-bind-nohost-1.c                        | 105 ++++++++++++++++
 .../libgomp.oacc-c-c++-common/routine-nohost-1.c   |  18 +++
 .../testsuite/libgomp.oacc-fortran/routine-6.f90   |   2 +-
 30 files changed, 589 insertions(+), 165 deletions(-)

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  <tho...@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  <clt...@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  <tho...@codesourcery.com>
+
+       * c-parser.c (c_parser_oacc_clause_bind, c_parser_oacc_routine)
+       (c_finish_oacc_routine): Update.
+
 2015-11-12  Nathan Sidwell  <nat...@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  <tho...@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  <tho...@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  <tho...@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  <ce...@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  <tho...@codesourcery.com>
+
+       * lto.c (read_cgraph_and_symbols) [ACCEL_COMPILER]: Handle OpenACC
+       bind clauses.
+
 2015-08-31  Nathan Sidwell  <nat...@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  <tho...@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  <ce...@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  <tho...@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  <jnor...@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


Grüße
 Thomas

Reply via email to