diff mbox

[gomp4] OpenACC routine directive

Message ID 541777A6.7090908@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis Sept. 15, 2014, 11:35 p.m. UTC 
This patch adds initial support for the OpenACC routine directive. It's
not complete just yet because it doesn't implement any of the optional
clauses, except for the optional function/subroutine name. As such, it
doesn't go beyond marking functions with the "omp declare target" attribute.

According to the OpenACC technical committee, the routine clause will be
revised in the next OpenACC 2.x release. In particular, function
definitions must have an 'acc routine' associated with it. My
understanding is 'acc routine' should also be visible at the call site,
but if it's not the compiler can treat it as a regular function call.
Furthermore, I've been told that it's not sufficient to place the
routine directive in an interface block by itself. E.g.

  interface
    recursive function fact (x)
      !$acc routine
      integer, intent(in) :: x
      integer :: fact
    end function fact
  end interface
  integer, parameter :: n = 10
  integer :: a(n), i
  !$acc parallel
  !$acc loop
  do i = 1, n
     a(i) = fact (i)
  end do
  !$acc end parallel
  do i = 1, n
     write (*, "(I10)") a(i)
  end do
end
recursive function fact (x) result (res)
  integer, intent(in) :: x
  integer :: res
  if (x < 1) then
     res = 1
  else
     res = x * fact (x - 1)
  end if
end function fact

This will result in a runtime failure because gcc will not generate an
accelerated version of fact. The justification for this is that fortran
lacks a file scope, so 'acc routine' wouldn't be visible to fact.

Is this patch OK for gomp-4_0-branch?

Thanks,
Cesar

Comments

Tobias Burnus Sept. 16, 2014, 7:24 a.m. UTC | #1 
Cesar Philippidis wrote:
> This patch adds initial support for the OpenACC routine directive. It's
> not complete just yet because it doesn't implement any of the optional
> clauses, except for the optional function/subroutine name. As such, it
> doesn't go beyond marking functions with the "omp declare target" attribute.

For the Fortran side: As you currently use the OpenMP implementation, it
should work, but if you later add support for clauses, recall that you may
need to store those also in the .mod files (cf. module.c). (That's only needed
if information from the clauses has to be propagated to the the call site.)


> My understanding is 'acc routine' should also be visible at the call site,
> but if it's not the compiler can treat it as a regular function call.

... which means that it has to reject it, unless the compiler can (e.g. via
LTO) inline the function or find out that the function also exist on the
accelerator.

> Furthermore, I've been told that it's not sufficient to place the
> routine directive in an interface block by itself. 

Well, that's also not different to C/C++: If you just use "#pragma acc routine"
in the header file, where you declare the function, it also won't work when
you don't have the pragma for the definition.

> The justification for this is that fortran
> lacks a file scope, so 'acc routine' wouldn't be visible to fact.

True. However,in modern Fortran use, you would use not use an interface
block (unless the function is written in, e.g., C). But you would use
a module - or an internal procedure (nested procedure with "contains").

If you want, you can also check that an omp declare target/acc routine
in an interface block matches the one in a subroutine/function by
adding a check in resolve.c's resolve_global_procedure.


> Is this patch OK for gomp-4_0-branch?

The Fortran part looks good to me.

Tobias

> 2014-09-15  Cesar Philippidis  <cesar@codesourcery.com>
>	gcc/fortran/
>	* gfortran.h (ST_OACC_ROUTINE): New statement enum.
>	* match.h (gfc_match_oacc_routine): New prototype.
>	* openmp.c (gfc_match_oacc_routine): New function.
>	* parse.c (decode_oacc_directive): Handle the routine directive.
>	(next_statement): Handle ST_OACC_ROUTINE.
>	(gfc_ascii_statement): Likewise.
Thomas Schwinge Sept. 23, 2014, 1:05 p.m. UTC | #2 
Hi Cesar!

On Mon, 15 Sep 2014 16:35:02 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> This patch adds initial support for the OpenACC routine directive.

Thanks!

> It's
> not complete just yet because it doesn't implement any of the optional
> clauses, except for the optional function/subroutine name. As such, it
> doesn't go beyond marking functions with the "omp declare target" attribute.

Reviewing the C front end bits.

I wonder if we can integrate this better (maybe as a follow-up patch),
for example, inside an OpenMP »#pragma omp declare target« region,
current_omp_declare_target_attribute is set to non-zero, and can then
conveniently be used in gcc/c/c-decl.c:c_decl_attributes.  (Of course,
OpenACC's »#pragma acc routine« is a bit different; doesn't have a »end«
tag.)  I have not yet figured out whether we can do something similar,
for example to avoid all the c_parser_declaration_or_fndef call-sites
changes.

> --- a/gcc/c/c-parser.c
> +++ b/gcc/c/c-parser.c
> @@ -1678,6 +1687,9 @@ c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
|  	  if (omp_declare_simd_clauses.exists ()
>  	      || !vec_safe_is_empty (parser->cilk_simd_fn_tokens))
>  	    c_finish_omp_declare_simd (parser, NULL_TREE, NULL_TREE,
>  				       omp_declare_simd_clauses);
> +	  else
> +	    c_finish_oacc_routine (parser, NULL_TREE,
> +				      oacc_routine_clauses, oacc_routine_named);
>  	  c_parser_skip_to_end_of_block_or_statement (parser);
>  	  return;
>  	}

Please be more explicit.  Using such "catch-all else statements", the
next reader of this code will wonder: why is
»!(omp_declare_simd_clauses.exists () || !vec_safe_is_empty
(parser->cilk_simd_fn_tokens))« the condition for calling
c_finish_oacc_routine?  (Likewise in other places.)

> +#define OACC_ROUTINE_CLAUSE_MASK					\
> +	(OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PROC_BIND)

PRAGMA_OMP_CLAUSE_NONE (for the moment), not PRAGMA_OMP_CLAUSE_PROC_BIND.

We're missing test cases for the different syntax options that are
described in the OpenACC specification, especially the invalid ones.  For
example, should gcc/testsuite/c-c++-common/goacc/pragma_context.c be
extended?


Grüße,
 Thomas
diff mbox

Patch

2014-09-15  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/c-family/
	* c-pragma.c (oacc_pragmas): Add the "routine" pragma.
	* c-pragma.h (pragma_kind): Add PRAGMA_OACC_ROUTINE.

	gcc/c/
	* c-parser.c (struct c_parser): New oacc_routines member.
	(c_parser_external_declaration): Update call to
	c_parser_declaration_or_fndef.
	(c_parser_declaration_or_fndef): Add oacc_routine_clauses and
	oacc_routine_named parameters. Use them in calls to
	c_finish_oacc_routine and c_parser_declaration_or_fndef.
	(c_parser_compound_statement_nostart): Update call to
	c_parser_declaration_or_fndef.
	(c_parser_label): Likewise.
	(c_parser_for_statement): Likewise.
	(c_parser_objc_methodprotolist): Likewise.
	(c_parser_pragma): Handle PRAGMA_OACC_ROUTINE.
	(OACC_ROUTINE_CLAUSE_MASK): New macro. 
	(c_parser_oacc_routine): New function.
	(c_finish_oacc_routine): New function.
	(c_parser_omp_for_loop): Update calls to c_parser_declaration_or_fndef.
	(c_parser_omp_declare_simd): Likewise.
	(c_parse_file): Initialize oacc_routines.

	gcc/fortran/
	* gfortran.h (ST_OACC_ROUTINE): New statement enum.
	* match.h (gfc_match_oacc_routine): New prototype.
	* openmp.c (gfc_match_oacc_routine): New function.
	* parse.c (decode_oacc_directive): Handle the routine directive.
	(next_statement): Handle ST_OACC_ROUTINE.
	(gfc_ascii_statement): Likewise.

	gcc/testsuite/
	* c-c++-common/goacc/routine-1.c: New test.
	* c-c++-common/goacc/routine-2.c: New test.
	* gfortran.dg/goacc/routine-1.f90: New test.
	* gfortran.dg/goacc/routine-2.f90: New test.


diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c
index 2d9071a..2ff99f5 100644
--- a/gcc/c-family/c-pragma.c
+++ b/gcc/c-family/c-pragma.c
@@ -1177,6 +1177,7 @@  static const struct omp_pragma_def oacc_pragmas[] = {
   { "kernels", PRAGMA_OACC_KERNELS },
   { "loop", PRAGMA_OACC_LOOP },
   { "parallel", PRAGMA_OACC_PARALLEL },
+  { "routine", PRAGMA_OACC_ROUTINE },
   { "update", PRAGMA_OACC_UPDATE },
   { "wait", PRAGMA_OACC_WAIT },
 };
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 4722d51..c097857 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -32,6 +32,7 @@  typedef enum pragma_kind {
   PRAGMA_OACC_KERNELS,
   PRAGMA_OACC_LOOP,
   PRAGMA_OACC_PARALLEL,
+  PRAGMA_OACC_ROUTINE,
   PRAGMA_OACC_UPDATE,
   PRAGMA_OACC_WAIT,
   PRAGMA_OMP_ATOMIC,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 09df223..0632c69 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -215,6 +215,10 @@  typedef struct GTY(()) c_parser {
   /* Buffer to hold all the tokens from parsing the vector attribute for the
      SIMD-enabled functions (formerly known as elemental functions).  */
   vec <c_token, va_gc> *cilk_simd_fn_tokens;
+
+  /* OpenACC specific parser information.  */
+
+  vec <tree, va_gc> *oacc_routines;
 } c_parser;
 
 
@@ -1150,7 +1154,8 @@  enum c_parser_prec {
 static void c_parser_external_declaration (c_parser *);
 static void c_parser_asm_definition (c_parser *);
 static void c_parser_declaration_or_fndef (c_parser *, bool, bool, bool,
-					   bool, bool, tree *, vec<c_token>);
+					   bool, bool, tree *, vec<c_token>,
+					   tree, bool);
 static void c_parser_static_assert_declaration_no_semi (c_parser *);
 static void c_parser_static_assert_declaration (c_parser *);
 static void c_parser_declspecs (c_parser *, struct c_declspecs *, bool, bool,
@@ -1238,6 +1243,8 @@  static bool c_parser_pragma (c_parser *, enum pragma_context);
 static bool c_parser_omp_target (c_parser *, enum pragma_context);
 static void c_parser_omp_end_declare_target (c_parser *);
 static void c_parser_omp_declare (c_parser *, enum pragma_context);
+static void c_parser_oacc_routine (c_parser *parser, enum pragma_context
+				      context);
 
 /* These Objective-C parser functions are only ever called when
    compiling Objective-C.  */
@@ -1417,12 +1424,13 @@  c_parser_external_declaration (c_parser *parser)
 	 only tell which after parsing the declaration specifiers, if
 	 any, and the first declarator.  */
       c_parser_declaration_or_fndef (parser, true, true, true, false, true,
-				     NULL, vNULL);
+				     NULL, vNULL, NULL_TREE, false);
       break;
     }
 }
 
 static void c_finish_omp_declare_simd (c_parser *, tree, tree, vec<c_token>);
+static void c_finish_oacc_routine (c_parser *, tree, tree, bool);
 
 /* Parse a declaration or function definition (C90 6.5, 6.7.1, C99
    6.7, 6.9.1).  If FNDEF_OK is true, a function definition is
@@ -1500,7 +1508,8 @@  c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
 			       bool static_assert_ok, bool empty_ok,
 			       bool nested, bool start_attr_ok,
 			       tree *objc_foreach_object_declaration,
-			       vec<c_token> omp_declare_simd_clauses)
+			       vec<c_token> omp_declare_simd_clauses,
+			       tree oacc_routine_clauses, bool oacc_routine_named)
 {
   struct c_declspecs *specs;
   tree prefix_attrs;
@@ -1678,6 +1687,9 @@  c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
 	      || !vec_safe_is_empty (parser->cilk_simd_fn_tokens))
 	    c_finish_omp_declare_simd (parser, NULL_TREE, NULL_TREE,
 				       omp_declare_simd_clauses);
+	  else
+	    c_finish_oacc_routine (parser, NULL_TREE,
+				      oacc_routine_clauses, oacc_routine_named);
 	  c_parser_skip_to_end_of_block_or_statement (parser);
 	  return;
 	}
@@ -1774,6 +1786,9 @@  c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
 		      || !vec_safe_is_empty (parser->cilk_simd_fn_tokens))
 		    c_finish_omp_declare_simd (parser, d, NULL_TREE,
 					       omp_declare_simd_clauses);
+		  else
+		    c_finish_oacc_routine (parser, d, oacc_routine_clauses,
+					      oacc_routine_named);
 		}
 	      else
 		{
@@ -1787,6 +1802,10 @@  c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
 		      || !vec_safe_is_empty (parser->cilk_simd_fn_tokens))
 		    c_finish_omp_declare_simd (parser, d, NULL_TREE,
 					       omp_declare_simd_clauses);
+		  else
+		    c_finish_oacc_routine (parser, d, oacc_routine_clauses,
+					      oacc_routine_named);
+
 		  start_init (d, asm_name, global_bindings_p ());
 		  init_loc = c_parser_peek_token (parser)->location;
 		  init = c_parser_initializer (parser);
@@ -1832,6 +1851,9 @@  c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
 		    temp_store_parm_decls (d, parms);
 		  c_finish_omp_declare_simd (parser, d, parms,
 					     omp_declare_simd_clauses);
+		  c_finish_oacc_routine (parser, d, oacc_routine_clauses,
+					    oacc_routine_named);
+
 		  if (parms)
 		    temp_pop_parm_decls ();
 		}
@@ -1938,13 +1960,17 @@  c_parser_declaration_or_fndef (c_parser *parser, bool fndef_ok,
 	 function definitions either.  */
       while (c_parser_next_token_is_not (parser, CPP_EOF)
 	     && c_parser_next_token_is_not (parser, CPP_OPEN_BRACE))
-	c_parser_declaration_or_fndef (parser, false, false, false,
-				       true, false, NULL, vNULL);
+	c_parser_declaration_or_fndef (parser, false, false, false, true,
+				       false, NULL, vNULL, NULL_TREE, false);
       store_parm_decls ();
       if (omp_declare_simd_clauses.exists ()
 	  || !vec_safe_is_empty (parser->cilk_simd_fn_tokens))
 	c_finish_omp_declare_simd (parser, current_function_decl, NULL_TREE,
 				   omp_declare_simd_clauses);
+      else
+	c_finish_oacc_routine (parser, current_function_decl,
+				  oacc_routine_clauses, oacc_routine_named);
+
       DECL_STRUCT_FUNCTION (current_function_decl)->function_start_locus
 	= c_parser_peek_token (parser)->location;
       fnbody = c_parser_compound_statement (parser);
@@ -4585,7 +4611,7 @@  c_parser_compound_statement_nostart (c_parser *parser)
 	  last_label = false;
 	  mark_valid_location_for_stdc_pragma (false);
 	  c_parser_declaration_or_fndef (parser, true, true, true, true,
-					 true, NULL, vNULL);
+					 true, NULL, vNULL, NULL_TREE, false);
 	  if (last_stmt)
 	    pedwarn_c90 (loc, OPT_Wdeclaration_after_statement,
 			 "ISO C90 forbids mixed declarations and code");
@@ -4610,7 +4636,8 @@  c_parser_compound_statement_nostart (c_parser *parser)
 	      last_label = false;
 	      mark_valid_location_for_stdc_pragma (false);
 	      c_parser_declaration_or_fndef (parser, true, true, true, true,
-					     true, NULL, vNULL);
+					     true, NULL, vNULL, NULL_TREE,
+					     false);
 	      /* Following the old parser, __extension__ does not
 		 disable this diagnostic.  */
 	      restore_extension_diagnostics (ext);
@@ -4747,7 +4774,7 @@  c_parser_label (c_parser *parser)
 					 /*static_assert_ok*/ true,
 					 /*empty_ok*/ true, /*nested*/ true,
 					 /*start_attr_ok*/ true, NULL,
-					 vNULL);
+					 vNULL, NULL_TREE, false);
 	}
     }
 }
@@ -5451,7 +5478,8 @@  c_parser_for_statement (c_parser *parser, bool ivdep)
       else if (c_parser_next_tokens_start_declaration (parser))
 	{
 	  c_parser_declaration_or_fndef (parser, true, true, true, true, true, 
-					 &object_expression, vNULL);
+					 &object_expression, vNULL, NULL_TREE,
+					 false);
 	  parser->objc_could_be_foreach_context = false;
 	  
 	  if (c_parser_next_token_is_keyword (parser, RID_IN))
@@ -5480,7 +5508,8 @@  c_parser_for_statement (c_parser *parser, bool ivdep)
 	      ext = disable_extension_diagnostics ();
 	      c_parser_consume_token (parser);
 	      c_parser_declaration_or_fndef (parser, true, true, true, true,
-					     true, &object_expression, vNULL);
+					     true, &object_expression, vNULL,
+					     NULL_TREE, false);
 	      parser->objc_could_be_foreach_context = false;
 	      
 	      restore_extension_diagnostics (ext);
@@ -8568,8 +8597,9 @@  c_parser_objc_methodprotolist (c_parser *parser)
 	      c_parser_consume_token (parser);
 	    }
 	  else
-	    c_parser_declaration_or_fndef (parser, false, false, true,
-					   false, true, NULL, vNULL);
+	    c_parser_declaration_or_fndef (parser, false, false, true,false,
+					   true, NULL, vNULL, NULL_TREE,
+					   false);
 	  break;
 	}
     }
@@ -9515,6 +9545,10 @@  c_parser_pragma (c_parser *parser, enum pragma_context context)
 
   switch (id)
     {
+    case PRAGMA_OACC_ROUTINE:
+      c_parser_oacc_routine (parser, context);
+      return false;
+
     case PRAGMA_OACC_UPDATE:
       if (context != pragma_compound)
 	{
@@ -12073,6 +12107,161 @@  c_parser_oacc_parallel (location_t loc, c_parser *parser, char *p_name)
   return stmt;
 }
 
+/* OpenACC 2.0: FIXME
+   # pragma acc routine oacc-routine-clause[optseq] new-line
+     function-definition
+
+   # pragma acc routine ( name ) oacc-routine-clause[optseq] new-line
+
+   LOC is the location of the #pragma token.
+*/
+
+#define OACC_ROUTINE_CLAUSE_MASK					\
+	(OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PROC_BIND)
+
+static void
+c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
+{
+  tree name = NULL_TREE;
+  location_t here = c_parser_peek_token (parser)->location;
+
+  c_parser_consume_pragma (parser);
+
+  /* Scan for optional '( name )'.  */
+  if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN)
+    {
+      c_parser_consume_token (parser);
+
+      if (c_parser_next_token_is_not (parser, CPP_NAME)
+	  || c_parser_peek_token (parser)->id_kind != C_ID_ID)
+	c_parser_error (parser, "expected identifier");
+
+      // name should be an IDENTIFIER_NODE
+      name = c_parser_peek_token (parser)->value;
+
+      if (name == NULL_TREE)
+	{
+	  undeclared_variable (c_parser_peek_token (parser)->location,
+			       c_parser_peek_token (parser)->value);
+	  name = error_mark_node;
+	}
+
+      c_parser_consume_token (parser);
+
+      if (name == error_mark_node)
+	return;
+
+      c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, 0);
+    }
+
+  /* Build a chain of clauses.  */
+  parser->in_pragma = true;
+  tree clauses = NULL_TREE;
+  clauses = c_parser_omp_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
+				"#pragma acc routine");
+
+  /* Check of the presence if gang, worker, vector and seq clauses, and
+     throw an error if more than one of those clauses is specified.  */
+  int parallelism = 0;
+  tree c;
+
+  for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    switch (OMP_CLAUSE_CODE (c))
+      {
+      case OMP_CLAUSE_GANG:
+      case OMP_CLAUSE_WORKER:
+      case OMP_CLAUSE_VECTOR:
+	/* FIXME: enable seq.
+	   case OMP_CLAUSE_SEQ: */
+	++parallelism;
+	break;
+      default:
+	break;
+      }
+
+  if (parallelism > 1)
+    {
+      error_at (here, "invalid combination of gang, worker, vector or seq for"
+		"%<#pragma acc routine%>");
+    }
+
+  if (name)
+    {
+      TREE_CHAIN (name) = clauses;
+      vec_safe_push (parser->oacc_routines, name);
+    }
+  else
+    {
+      if (context != pragma_external)
+	{
+	  c_parser_error (parser, "%<#pragma acc routine%> must be "
+			  "followed by function declaration or definition");
+	  return;
+	}
+
+      if (c_parser_next_token_is (parser, CPP_KEYWORD)
+	  && c_parser_peek_token (parser)->keyword == RID_EXTENSION)
+	{
+	  int ext = disable_extension_diagnostics ();
+	  do
+	    c_parser_consume_token (parser);
+	  while (c_parser_next_token_is (parser, CPP_KEYWORD)
+		 && c_parser_peek_token (parser)->keyword
+		 == RID_EXTENSION);
+	  c_parser_declaration_or_fndef (parser, true, true, true, false,
+					 true, NULL, vNULL, clauses, true);
+	  restore_extension_diagnostics (ext);
+	}
+      else
+	c_parser_declaration_or_fndef (parser, true, true, true, false,
+				       true, NULL, vNULL, clauses, true);
+    }
+}
+
+static void
+c_finish_oacc_routine (c_parser *parser, tree fndecl, tree clauses,
+			  bool named)
+{
+  if (fndecl == NULL_TREE || TREE_CODE (fndecl) != FUNCTION_DECL)
+    {
+      if (!named)
+	return;
+
+      error ("%<#pragma acc routine%> not immediately followed by "
+	     "a function declaration or definition");
+      gcc_unreachable();
+      return;
+    }
+
+  if (!named)
+    {
+      bool found = false;
+      int i;
+      tree t;
+
+      for (i = 0; vec_safe_iterate (parser->oacc_routines, i, &t); i++)
+	{
+	  if (!strcmp (IDENTIFIER_POINTER (DECL_NAME (fndecl)),
+		       IDENTIFIER_POINTER (t)))
+	    {
+	      found = true;
+	      clauses = TREE_CHAIN (t);
+	      break;
+	    }
+	}
+
+      if (!found)
+	return;
+    }
+
+  if (clauses != NULL_TREE)
+    clauses = tree_cons (NULL_TREE, clauses, NULL_TREE);
+  clauses = build_tree_list (get_identifier ("omp declare target"),
+			     clauses);
+  TREE_CHAIN (clauses) = DECL_ATTRIBUTES (fndecl);
+  DECL_ATTRIBUTES (fndecl) = clauses;
+}
+
 /* OpenACC 2.0:
    # pragma acc update oacc-update-clause[optseq] new-line
 */
@@ -12665,7 +12854,7 @@  c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
 	  if (i > 0)
 	    vec_safe_push (for_block, c_begin_compound_stmt (true));
 	  c_parser_declaration_or_fndef (parser, true, true, true, true, true,
-					 NULL, vNULL);
+					 NULL, vNULL, NULL_TREE, false);
 	  decl = check_for_loop_decls (for_loc, flag_isoc99);
 	  if (decl == NULL)
 	    goto error_init;
@@ -13854,12 +14043,12 @@  c_parser_omp_declare_simd (c_parser *parser, enum pragma_context context)
 	  while (c_parser_next_token_is (parser, CPP_KEYWORD)
 		 && c_parser_peek_token (parser)->keyword == RID_EXTENSION);
 	  c_parser_declaration_or_fndef (parser, true, true, true, false, true,
-					 NULL, clauses);
+					 NULL, clauses, NULL_TREE, false);
 	  restore_extension_diagnostics (ext);
 	}
       else
 	c_parser_declaration_or_fndef (parser, true, true, true, false, true,
-				       NULL, clauses);
+				       NULL, clauses, NULL_TREE, false);
       break;
     case pragma_struct:
     case pragma_param:
@@ -13879,7 +14068,8 @@  c_parser_omp_declare_simd (c_parser *parser, enum pragma_context context)
 	  if (c_parser_next_tokens_start_declaration (parser))
 	    {
 	      c_parser_declaration_or_fndef (parser, true, true, true, true,
-					     true, NULL, clauses);
+					     true, NULL, clauses, NULL_TREE,
+					     false);
 	      restore_extension_diagnostics (ext);
 	      break;
 	    }
@@ -13888,7 +14078,7 @@  c_parser_omp_declare_simd (c_parser *parser, enum pragma_context context)
       else if (c_parser_next_tokens_start_declaration (parser))
 	{
 	  c_parser_declaration_or_fndef (parser, true, true, true, true, true,
-					 NULL, clauses);
+					 NULL, clauses, NULL_TREE, false);
 	  break;
 	}
       c_parser_error (parser, "%<#pragma omp declare simd%> must be followed by "
@@ -15230,6 +15420,8 @@  c_parse_file (void)
   if (tparser.tokens == &tparser.tokens_buf[0])
     the_parser->tokens = &the_parser->tokens_buf[0];
 
+  the_parser->oacc_routines = NULL;
+
   /* Initialize EH, if we've been told to do so.  */
   if (flag_exceptions)
     using_eh_for_cleanups ();
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 755e62e..a379946 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -221,7 +221,7 @@  typedef enum
   ST_OACC_END_DATA, ST_OACC_HOST_DATA, ST_OACC_END_HOST_DATA, ST_OACC_LOOP,
   ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE, ST_OACC_WAIT,
   ST_OACC_CACHE, ST_OACC_KERNELS_LOOP, ST_OACC_END_KERNELS_LOOP,
-  ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA,
+  ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA, ST_OACC_ROUTINE,
   ST_OMP_ATOMIC, ST_OMP_BARRIER, ST_OMP_CRITICAL, ST_OMP_END_ATOMIC,
   ST_OMP_END_CRITICAL, ST_OMP_END_DO, ST_OMP_END_MASTER, ST_OMP_END_ORDERED,
   ST_OMP_END_PARALLEL, ST_OMP_END_PARALLEL_DO, ST_OMP_END_PARALLEL_SECTIONS,
diff --git a/gcc/fortran/match.h b/gcc/fortran/match.h
index 3624638..11e199d 100644
--- a/gcc/fortran/match.h
+++ b/gcc/fortran/match.h
@@ -136,6 +136,7 @@  match gfc_match_oacc_parallel (void);
 match gfc_match_oacc_parallel_loop (void);
 match gfc_match_oacc_enter_data (void);
 match gfc_match_oacc_exit_data (void);
+match gfc_match_oacc_routine (void);
 
 /* OpenMP directive matchers.  */
 match gfc_match_omp_eos (void);
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index ecb20d2..bb3b1b7 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -1367,6 +1367,79 @@  gfc_match_oacc_cache (void)
 }
 
 
+match
+gfc_match_oacc_routine (void)
+{
+  locus old_loc;
+  gfc_symbol *sym;
+  match m;
+
+  old_loc = gfc_current_locus;
+
+  m = gfc_match (" (");
+
+  if (gfc_current_ns->proc_name
+      && gfc_current_ns->proc_name->attr.if_source == IFSRC_IFBODY
+      && m == MATCH_YES)
+    {
+      gfc_error ("Only the !$ACC ROUTINE form without "
+		 "list is allowed in interface block at %C");
+      goto cleanup;
+    }
+
+  if (m == MATCH_NO
+      && gfc_current_ns->proc_name
+      && gfc_match_omp_eos () == MATCH_YES)
+    {
+      if (!gfc_add_omp_declare_target (&gfc_current_ns->proc_name->attr,
+				       gfc_current_ns->proc_name->name,
+				       &old_loc))
+	goto cleanup;
+      return MATCH_YES;
+    }
+
+  if (m != MATCH_YES)
+    return m;
+
+  /* Scan for a function name.  */
+  m = gfc_match_symbol (&sym, 0);
+
+  if (m != MATCH_YES)
+    {
+      gfc_error ("Syntax error in !$ACC ROUTINE ( NAME ) at %C");
+      gfc_current_locus = old_loc;
+      return MATCH_ERROR;
+    }
+
+  if (!sym->attr.external && !sym->attr.function && !sym->attr.subroutine)
+    {
+      gfc_error ("Syntax error in !$ACC ROUTINE ( NAME ) at %C, invalid"
+		 " function name '%s'", sym->name);
+      gfc_current_locus = old_loc;
+      return MATCH_ERROR;
+    }
+
+  if (gfc_match_char (')') != MATCH_YES)
+    {
+      gfc_error ("Syntax error in !$ACC ROUTINE ( NAME ) at %C, expecting"
+		 " ')' after NAME");
+      gfc_current_locus = old_loc;
+      return MATCH_ERROR;
+    }
+
+  if (gfc_match_omp_eos () != MATCH_YES)
+    {
+      gfc_error ("Unexpected junk after !$ACC ROUTINE at %C");
+      goto cleanup;
+    }
+  return MATCH_YES;
+
+cleanup:
+  gfc_current_locus = old_loc;
+  return MATCH_ERROR;
+}
+
+
 #define OMP_PARALLEL_CLAUSES \
   (OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_SHARED	\
    | OMP_CLAUSE_COPYIN | OMP_CLAUSE_REDUCTION | OMP_CLAUSE_IF		\
diff --git a/gcc/fortran/parse.c b/gcc/fortran/parse.c
index b2241d3..411dc27 100644
--- a/gcc/fortran/parse.c
+++ b/gcc/fortran/parse.c
@@ -644,6 +644,9 @@  decode_oacc_directive (void)
     case 'l':
       match ("loop", gfc_match_oacc_loop, ST_OACC_LOOP);
       break;
+    case 'r':
+      match ("routine", gfc_match_oacc_routine, ST_OACC_ROUTINE);
+      break;
     case 'u':
       match ("update", gfc_match_oacc_update, ST_OACC_UPDATE);
       break;
@@ -1356,7 +1359,7 @@  next_statement (void)
   case ST_EQUIVALENCE: case ST_NAMELIST: case ST_STATEMENT_FUNCTION: \
   case ST_TYPE: case ST_INTERFACE: case ST_OMP_THREADPRIVATE: \
   case ST_PROCEDURE: case ST_OMP_DECLARE_SIMD: case ST_OMP_DECLARE_REDUCTION: \
-  case ST_OMP_DECLARE_TARGET
+  case ST_OMP_DECLARE_TARGET: case ST_OACC_ROUTINE
 
 /* Block end statements.  Errors associated with interchanging these
    are detected in gfc_match_end().  */
@@ -1903,6 +1906,9 @@  gfc_ascii_statement (gfc_statement st)
     case ST_OACC_EXIT_DATA:
       p = "!$ACC EXIT DATA";
       break;
+    case ST_OACC_ROUTINE:
+      p = "!$ACC ROUTINE";
+      break;
     case ST_OMP_ATOMIC:
       p = "!$OMP ATOMIC";
       break;
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-1.c b/gcc/testsuite/c-c++-common/goacc/routine-1.c
new file mode 100644
index 0000000..1f89fdb
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -0,0 +1,35 @@ 
+void *malloc (__SIZE_TYPE__);
+void free (void *);
+
+#pragma acc routine
+int
+fact (int n)
+{
+  if (n == 0 || n == 1)
+    return 1;
+
+  return n * fact (n - 1);
+}
+
+int
+main(int argc, char **argv)
+{
+  int *a, i, n = 10;
+
+  a = (int *)malloc (sizeof (int) * n);
+
+#pragma acc parallel copy (a[0:n]) vector_length (5)
+  {
+#pragma acc loop
+    for (i = 0; i < n; i++)
+      a[i] = fact (i);
+  }
+
+  for (i = 0; i < n; i++)
+    if (fact (i) != a[i])
+      return -1;
+
+  free (a);
+
+  return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-2.c b/gcc/testsuite/c-c++-common/goacc/routine-2.c
new file mode 100644
index 0000000..fe2e7f7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-2.c
@@ -0,0 +1,36 @@ 
+void *malloc (__SIZE_TYPE__);
+void free (void *);
+
+#pragma acc routine (fact)
+
+int
+fact (int n)
+{
+  if (n == 0 || n == 1)
+    return 1;
+
+  return n * fact (n - 1);
+}
+
+int
+main(int argc, char **argv)
+{
+  int *a, i, n = 10;
+
+  a = (int *)malloc (sizeof (int) * n);
+
+#pragma acc parallel copy (a[0:n]) vector_length (5)
+  {
+#pragma acc loop
+    for (i = 0; i < n; i++)
+      a[i] = fact (i);
+  }
+
+  for (i = 0; i < n; i++)
+    if (fact (i) != a[i])
+      return -1;
+
+  free (a);
+
+  return 0;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-1.f90
new file mode 100644
index 0000000..67c5f11
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-1.f90
@@ -0,0 +1,37 @@ 
+! { dg-do compile }
+
+  integer, parameter :: n = 10
+  integer :: a(n), i
+  integer, external :: fact
+  i = 1
+  !$acc routine (fact)  ! { dg-error "Unexpected \\\!\\\$ACC ROUTINE" }
+  !$acc routine ()  ! { dg-error "Syntax error in \\\!\\\$ACC ROUTINE \\\( NAME \\\)" }
+  !$acc parallel
+  !$acc loop
+  do i = 1, n
+     a(i) = fact (i)
+     call incr (a(i))
+  end do
+  !$acc end parallel
+  do i = 1, n
+     write (*, "(I10)") a(i)
+  end do
+end
+recursive function fact (x) result (res)
+  integer, intent(in) :: x
+  integer :: res
+  res = 1
+  !$acc routine  ! { dg-error "Unexpected \\\!\\\$ACC ROUTINE" }
+  if (x < 1) then
+     res = 1
+  else
+     res = x * fact (x - 1)
+  end if
+end function fact
+subroutine incr (x)
+  integer, intent(inout) :: x
+  integer i
+  i = 0
+  !$acc routine  ! { dg-error "Unexpected \\\!\\\$ACC ROUTINE" }
+  x = x + 1
+end subroutine incr
diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-2.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-2.f90
new file mode 100644
index 0000000..3be3351
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-2.f90
@@ -0,0 +1,17 @@ 
+! { dg-do compile }
+
+  module m1
+    contains
+    recursive function mfact (x) result (res)
+      integer, intent(in) :: x
+      integer :: res
+      integer i
+      i = 0
+      !$acc routine  ! { dg-error "Unexpected \\\!\\\$ACC ROUTINE" }
+      if (x < 1) then
+         res = 1
+      else
+         res = x * mfact (x - 1)
+      end if
+    end function mfact
+  end module m1