diff mbox

[gomp4] Middle end bits for routines

Message ID 55CF58E5.2090001@acm.org
State New
Headers show

Commit Message

Nathan Sidwell Aug. 15, 2015, 3:21 p.m. UTC
I've committed this to the gomp4 branch.  It extends the  'oacc function' 
attribute's dimension handling to deal with routines.  With the latter we now 
use TREE_PURPOSE of the dimension list to indicate whether the routine may or 
may not spawn a partitioned loop on a particular axis.  I had to tweak the 
validate_dims hook to tell it what the outermost such axis is (or -1 for a 
non-routine).

The patch was complicated y the rather baroque handling of the routine directive 
in  the C and C++ frontends.  Now I see that handling is based on a 
misinterpretation of the specification.  I'll clean that up shortly.

nathan
diff mbox

Patch

2015-08-15  Nathan Sidwell  <nathan@codesourcery.com>

	* config/nvptx/nvptx.c (nvptx_reorg): Examine TREE_PURPOSE of
	dimensions.
	(nvptx_record_offload_symbol): Adjust.
	(nvptx_validate_dims): Adjust.
	* omp-low.c (replace_oacc_fn_attrib): Detect if we're the first
	attribute.
	(set_oacc_fn_attrub): Swap FN and CLAUSE parameters for
	consistency.
	(build_oacc_routine_dims): New.
	(expand_omp_target): Adjust set_oacc_fn_attrib call.
	(execute_oacc_transform): Deal with routine dimensions.
	(default_goacc_validate_dims): Add FN_LEVEL parameter.
	* omp-low.h (replace_oacc_fn_attrib, build_oacc_routine_dims): Declare.
	* target.def (validate_dims): Add FN_LEVEL arg.
	* targhooks.h (default_goacc_validate_dims): Adjust.
	* doc/tm.texi: Rebuilt.

	testsuite/
	* c-c++-common/goacc/routine-4.c: Adjust expected error.  Delete
	bogus tests.

	fortran/
	*  f95-lang.c (gfc_attribs): oacc function attribute can take
	list.
	* gfortran.h (struct symbol_attribute): Add more bits to
	oacc_routine field.
	* openmp.c (gfc_oacc_routine_dims): New.
	(gfc_match_oacc_routins): Call it.
	* trans-decl.c: Include gomp-constants.h.
	(add_attributes_to_decl): Create oacc function dimension data.

	cp/
	* parser.c (cp_parser_oacc_routine_check_parallelism): Delete.
	(cp_parser_oacc_routine): Don't check parallelism here.
	(cp_parser_late_parssing_oacc_routine): Use build_oacc_routine_dims.

	c/
	* c-parser.c (c_parser_oacc_routine): Don't check loop axes here,
	use build_oacc_routine_dims.

	c-family/
	* c-common.c (c_common_attribs): oacc function can take list.
	
Index: gcc/c/c-parser.c
===================================================================
--- gcc/c/c-parser.c	(revision 226860)
+++ gcc/c/c-parser.c	(working copy)
@@ -13291,7 +13291,6 @@  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);
 
@@ -13329,30 +13328,6 @@  c_parser_oacc_routine (c_parser *parser,
 				       "#pragma acc routine",
 				       OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK);
 
-  /* 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:
-      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;
@@ -13422,14 +13397,16 @@  c_finish_oacc_routine (c_parser *parser,
 	return;
     }
 
+  /* Process for function attrib  */
+  tree dims = build_oacc_routine_dims (clauses);
+  replace_oacc_fn_attrib (fndecl, dims);
+
+  /* Also attach as a declare.  */
   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)
-    = tree_cons (get_identifier ("oacc function"),
-		 NULL_TREE, clauses);
+    = tree_cons (get_identifier ("omp declare target"),
+		 clauses, DECL_ATTRIBUTES (fndecl));
 }
 
 /* OpenACC 2.0:
Index: gcc/c-family/c-common.c
===================================================================
--- gcc/c-family/c-common.c	(revision 226860)
+++ gcc/c-family/c-common.c	(working copy)
@@ -823,7 +823,7 @@  const struct attribute_spec c_common_att
   { "bnd_instrument",         0, 0, true, false, false,
 			      handle_bnd_instrument, false },
   { "oacc declare",           0, -1, true,  false, false, NULL, false },
-  { "oacc function",          0, 0, true,  false, false, NULL, false },
+  { "oacc function",          0, -1, true,  false, false, NULL, false },
   { NULL,                     0, 0, false, false, false, NULL, false }
 };
 
Index: gcc/config/nvptx/nvptx.c
===================================================================
--- gcc/config/nvptx/nvptx.c	(revision 226860)
+++ gcc/config/nvptx/nvptx.c	(working copy)
@@ -3079,9 +3079,10 @@  nvptx_reorg (void)
 
       for (ix = 0; ix != GOMP_DIM_MAX; ix++, dims = TREE_CHAIN (dims))
 	{
-	  HOST_WIDE_INT size = TREE_INT_CST_LOW (TREE_VALUE (dims));
+	  int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
+	  tree allowed = TREE_PURPOSE (dims);
 
-	  if (size > 1 || (!size && !TREE_PURPOSE (dims)))
+	  if (size != 1 && !(allowed && integer_zerop (allowed)))
 	    mask |= GOMP_DIM_MASK (ix);
 	}
       /* If there is worker neutering, there must be vector
@@ -3188,10 +3189,10 @@  nvptx_record_offload_symbol (tree decl)
 
 	for (ix = 0; ix != GOMP_DIM_MAX; ix++, dims = TREE_CHAIN (dims))
 	  {
-	    HOST_WIDE_INT size = TREE_INT_CST_LOW (TREE_VALUE (dims));
+	    int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
 
 	    gcc_assert (!TREE_PURPOSE (dims));
-	    fprintf (asm_out_file, ", " HOST_WIDE_INT_PRINT_HEX, size);
+	    fprintf (asm_out_file, ", %#x", size);
 	  }
 
 	fprintf (asm_out_file, "\n");
@@ -3543,15 +3544,24 @@  nvptx_expand_builtin (tree exp, rtx targ
 #define PTX_VECTOR_LENGTH 32
 #define PTX_WORKER_LENGTH 32
 
-/* Validate compute dimensions, fill in non-unity defaults.  */
+/* Validate compute dimensions, fill in non-unity defaults.  FN_LEVEL
+   indicates the level at which a routine might spawn a loop.  It is
+   negative for non-routines.  */
 
 static bool
-nvptx_validate_dims (tree decl, int dims[])
+nvptx_validate_dims (tree decl, int dims[], int fn_level)
 {
   bool changed = false;
 
+  if (fn_level >= 0)
+    /* This is a routine.  All dimensions are dynamic and controlled
+       by the  calling function.  Because we permit a 1vx1wxNg
+       geometry, we can't take the opportunity to fix the vector
+       dimension inside a routine.  Perhaps we should?  */
+    return false;
+  
   /* If the worker size is not 1, the vector size must be 32.  If
-     the vector  size is not 1, it must be 32.  */
+     the vector size is not 1, it must be 32.  */
   if ((dims[GOMP_DIM_WORKER] > 1 || dims[GOMP_DIM_WORKER] == 0)
       || (dims[GOMP_DIM_VECTOR] > 1 || dims[GOMP_DIM_VECTOR] == 0))
     {
Index: gcc/cp/parser.c
===================================================================
--- gcc/cp/parser.c	(revision 226860)
+++ gcc/cp/parser.c	(working copy)
@@ -34368,34 +34368,6 @@  cp_parser_omp_declare (cp_parser *parser
   cp_parser_require_pragma_eol (parser, pragma_tok);
 }
 
-static void
-cp_parser_oacc_routine_check_parallelism (tree clauses, location_t loc)
-{
- /* 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:
-      case OMP_CLAUSE_SEQ:
-	++parallelism;
-	break;
-      default:
-	break;
-      }
-
-  if (parallelism > 1)
-    {
-      error_at (loc, "invalid combination of gang, worker, vector or seq for"
-		"%<#pragma acc routine%>");
-    }
-}
-
 /* OpenACC 2.0:
    # pragma acc routine oacc-routine-clause[optseq] new-line
      function-definition
@@ -34424,7 +34396,6 @@  cp_parser_oacc_routine (cp_parser *parse
 			enum pragma_context context)
 {
   tree name = NULL_TREE;
-  location_t here = cp_lexer_peek_token (parser->lexer)->location;
 
   //cp_lexer_consume_token (parser->lexer);
 
@@ -34466,8 +34437,6 @@  cp_parser_oacc_routine (cp_parser *parse
 					cp_lexer_peek_token (parser->lexer),
 					OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK);
 
-  cp_parser_oacc_routine_check_parallelism (clauses, here);
-
   TREE_CHAIN (name) = clauses;
   vec_safe_push (parser->named_oacc_routines, name);
 }
@@ -34481,7 +34450,6 @@  cp_parser_late_parsing_oacc_routine (cp_
   struct cp_token_cache *ce;
   cp_omp_declare_simd_data *data = parser->oacc_routine;
   int i;
-  location_t here = UNKNOWN_LOCATION;
 
   if (!data->error_seen && data->fndecl_seen)
     {
@@ -34499,7 +34467,6 @@  cp_parser_late_parsing_oacc_routine (cp_
     {
       cp_parser_push_lexer_for_tokens (parser, ce);
       parser->lexer->in_pragma = true;
-      here = cp_lexer_peek_token (parser->lexer)->location;
       gcc_assert (cp_lexer_peek_token (parser->lexer)->type == CPP_PRAGMA);
       cp_token *pragma_tok = cp_lexer_consume_token (parser->lexer);
       c = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
@@ -34519,13 +34486,13 @@  cp_parser_late_parsing_oacc_routine (cp_
 	}
     }
 
-  cp_parser_oacc_routine_check_parallelism (cl, here);
+  tree dims = build_oacc_routine_dims (cl);
+  attrs = tree_cons (get_identifier ("oacc function"), dims, attrs);
 
   if (cl != NULL_TREE)
     cl = tree_cons (NULL_TREE, cl, NULL_TREE);
 
-  attrs = build_tree_list (get_identifier ("omp declare target"), cl);
-  attrs = tree_cons (get_identifier ("oacc function"), NULL_TREE, attrs);
+  attrs = tree_cons (get_identifier ("omp declare target"), cl, attrs);
   data->fndecl_seen = true;
   return attrs;
 }
Index: gcc/doc/tm.texi
===================================================================
--- gcc/doc/tm.texi	(revision 226860)
+++ gcc/doc/tm.texi	(working copy)
@@ -5740,7 +5740,7 @@  usable.  In that case, the smaller the n
 to use it.
 @end deftypefn
 
-@deftypefn {Target Hook} bool TARGET_GOACC_VALIDATE_DIMS (tree, int @var{[]})
+@deftypefn {Target Hook} bool TARGET_GOACC_VALIDATE_DIMS (tree, int @var{[]}, @var{int})
 This hook should check the launch dimensions provided.  It should fill
 in anything that needs default to non-unity and verify non-defaults.
 Defaults are represented as -1.  Diagnostics should be issuedas 
Index: gcc/fortran/f95-lang.c
===================================================================
--- gcc/fortran/f95-lang.c	(revision 226860)
+++ gcc/fortran/f95-lang.c	(working copy)
@@ -106,7 +106,7 @@  static const struct attribute_spec gfc_a
     gfc_handle_omp_declare_target_attribute, false },
   { "oacc declare", 0, 0, true,  false, false,
     gfc_handle_omp_declare_target_attribute, false },
-  { "oacc function", 0, 0, true,  false, false,
+  { "oacc function", 0, -1, true,  false, false,
     gfc_handle_omp_declare_target_attribute, false },
   { NULL,		  0, 0, false, false, false, NULL, false }
 };
Index: gcc/fortran/gfortran.h
===================================================================
--- gcc/fortran/gfortran.h	(revision 226860)
+++ gcc/fortran/gfortran.h	(working copy)
@@ -875,8 +875,8 @@  typedef struct
   unsigned oacc_declare_device_resident:1;
   unsigned oacc_declare_link:1;
 
-  /* This is an OpenACC acclerator function.  */
-  unsigned oacc_function:1;
+  /* This is an OpenACC acclerator function at level N - 1  */
+  unsigned oacc_function:3;
 
   /* Attributes set by compiler extensions (!GCC$ ATTRIBUTES).  */
   unsigned ext_attr:EXT_ATTR_NUM;
Index: gcc/fortran/openmp.c
===================================================================
--- gcc/fortran/openmp.c	(revision 226860)
+++ gcc/fortran/openmp.c	(working copy)
@@ -1691,6 +1691,35 @@  gfc_match_oacc_cache (void)
   return MATCH_YES;
 }
 
+/* Determine the loop level for a routine.   */
+
+static int
+gfc_oacc_routine_dims (gfc_omp_clauses *clauses)
+{
+  int level = -1;
+
+  if (clauses)
+    {
+      unsigned mask = 0;
+
+      if (clauses->gang)
+	level = GOMP_DIM_GANG, mask |= GOMP_DIM_MASK (level);
+      if (clauses->worker)
+	level = GOMP_DIM_WORKER, mask |= GOMP_DIM_MASK (level);
+      if (clauses->vector)
+	level = GOMP_DIM_VECTOR, mask |= GOMP_DIM_MASK (level);
+      if (clauses->seq)
+	level = GOMP_DIM_MAX, mask |= GOMP_DIM_MASK (level);
+
+      if (mask != (mask & -mask))
+	gfc_error ("Multiple loop axes specified for routine");
+    }
+
+  if (level < 0)
+    level = GOMP_DIM_MAX;
+
+  return level;
+}
 
 match
 gfc_match_oacc_routine (void)
@@ -1743,6 +1772,12 @@  gfc_match_oacc_routine (void)
       }
     }
 
+  if (gfc_match_omp_eos () != MATCH_YES
+      && gfc_match_omp_clauses (&c, OACC_ROUTINE_CLAUSES,
+				OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK, false,
+				false, true) != MATCH_YES)
+    return MATCH_ERROR;
+
   if (sym != NULL)
     {
       n = gfc_get_oacc_routine_name ();
@@ -1760,20 +1795,12 @@  gfc_match_oacc_routine (void)
 				       gfc_current_ns->proc_name->name,
 				       &old_loc))
 	goto cleanup;
-      gfc_current_ns->proc_name->attr.oacc_function = 1;
+      gfc_current_ns->proc_name->attr.oacc_function
+	= gfc_oacc_routine_dims (c) + 1;
     }
   else
     gcc_unreachable ();
 
-  if (gfc_match_omp_eos () == MATCH_YES)
-    return MATCH_YES;
-
-  if (gfc_match_omp_clauses (&c, OACC_ROUTINE_CLAUSES,
-			     OACC_ROUTINE_CLAUSE_DEVICE_TYPE_MASK, false,
-			     false, true)
-      != MATCH_YES)
-    return MATCH_ERROR;
-
   if (n)
     n->clauses = c;
   else if (gfc_current_ns->oacc_routine)
Index: gcc/fortran/trans-decl.c
===================================================================
--- gcc/fortran/trans-decl.c	(revision 226860)
+++ gcc/fortran/trans-decl.c	(working copy)
@@ -49,6 +49,7 @@  along with GCC; see the file COPYING3.
 #include "trans-const.h"
 /* Only for gfc_trans_code.  Shouldn't need to include this.  */
 #include "trans-stmt.h"
+#include "gomp-constants.h"
 
 #define MAX_LABEL_VALUE 99999
 
@@ -1320,8 +1321,18 @@  add_attributes_to_decl (symbol_attribute
     }
 
   if (sym_attr.oacc_function)
-    list = tree_cons (get_identifier ("oacc function"),
-		      NULL_TREE, list);
+    {
+      tree dims = NULL_TREE;
+      int ix;
+      int level = sym_attr.oacc_function - 1;
+
+      for (ix = GOMP_DIM_MAX; ix--;)
+	dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
+			  integer_zero_node, dims);
+
+      list = tree_cons (get_identifier ("oacc function"),
+			dims, list);
+    }
 
   return list;
 }
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c	(revision 226860)
+++ gcc/omp-low.c	(working copy)
@@ -9319,13 +9319,17 @@  oacc_launch_pack (unsigned code, tree de
 void
 replace_oacc_fn_attrib (tree fn, tree dims)
 {
-  /* Simply cons onto the beginning of the list.  */
-  DECL_ATTRIBUTES (fn) =
-    tree_cons (get_identifier (OACC_FN_ATTRIB), dims, DECL_ATTRIBUTES (fn));
+  tree ident = get_identifier (OACC_FN_ATTRIB);
+  tree attribs = DECL_ATTRIBUTES (fn);
+
+  /* If we happen to be present as the first attrib, drop it.  */
+  if (attribs && TREE_PURPOSE (attribs) == ident)
+    attribs = TREE_CHAIN (attribs);
+  DECL_ATTRIBUTES (fn) = tree_cons (ident, dims, attribs);
 }
 
 static void
-set_oacc_fn_attrib (tree clauses, tree fn, vec<tree> *args)
+set_oacc_fn_attrib (tree fn, tree clauses, vec<tree> *args)
 {
   /* Must match GOMP_DIM ordering.  */
   static const omp_clause_code ids[] = 
@@ -9364,6 +9368,45 @@  set_oacc_fn_attrib (tree clauses, tree f
     }
 }
 
+/*  Process the routine's dimension clauess to generate an attribute
+    value.  Issue diagnostics as appropriate.  We default to SEQ
+    (OpenACC 2.5 clarifies this). All dimensions have a size of zero
+    (dynamic).  TREE_PURPOSE is set to indicate whether that dimension
+    can have a loop partitioned on it.  boolean_true_node indicates
+    yes, boolean_false_node indicates no.  */
+
+tree
+build_oacc_routine_dims (tree clauses)
+{
+  /* Must match GOMP_DIM ordering.  */
+  static const omp_clause_code ids[] = 
+    {OMP_CLAUSE_GANG, OMP_CLAUSE_WORKER, OMP_CLAUSE_VECTOR, OMP_CLAUSE_SEQ};
+  int ix;
+  int level = -1;
+
+  for (; clauses; clauses = OMP_CLAUSE_CHAIN (clauses))
+    for (ix = GOMP_DIM_MAX + 1; ix--;)
+      if (OMP_CLAUSE_CODE (clauses) == ids[ix])
+	{
+	  if (level >= 0)
+	    error_at (OMP_CLAUSE_LOCATION (clauses),
+		      "multiple loop axes specified for routine");
+	  level = ix;
+	  break;
+	}
+
+  if (level < 0)
+    level = GOMP_DIM_MAX;
+  
+  tree dims = NULL_TREE;
+
+  for (ix = GOMP_DIM_MAX; ix--;)
+    dims = tree_cons (build_int_cst (boolean_type_node, ix >= level),
+		      integer_zero_node, dims);
+
+  return dims;
+}
+
 tree
 get_oacc_fn_attrib (tree fn)
 {
@@ -9862,7 +9905,7 @@  expand_omp_target (struct omp_region *re
     case BUILT_IN_GOACC_PARALLEL:
       {
 	args.quick_push (gimple_omp_target_ganglocal_size (entry_stmt));
-	set_oacc_fn_attrib (clauses, child_fn, &args);
+	set_oacc_fn_attrib (child_fn, clauses, &args);
 	tagging = true;
       }
       /* FALLTHRU */
@@ -14624,6 +14667,7 @@  execute_oacc_transform ()
   basic_block bb;
   tree attrs = get_oacc_fn_attrib (current_function_decl);
   int dims[GOMP_DIM_MAX];
+  tree purpose[GOMP_DIM_MAX];
   
   if (!attrs)
     /* Not an offloaded function.  */
@@ -14632,36 +14676,47 @@  execute_oacc_transform ()
   {
     unsigned ix;
     tree pos = TREE_VALUE (attrs);
+    int fn_level = -1;
 
+    /* Make sure the attribute creator attached the dimension
+       information.  */
+    gcc_assert (pos);
+    
     for (ix = 0; ix != GOMP_DIM_MAX; ix++)
       {
-	if (!pos)
-	  dims[ix] = -1;
-	else
+	purpose[ix] = TREE_PURPOSE (pos);
+
+	if (purpose[ix])
 	  {
-	    tree val = TREE_VALUE (pos);
-	    
-	    dims[ix] = val ? TREE_INT_CST_LOW (val) : -2;
-	    pos = TREE_CHAIN (pos);
+	    if (purpose[ix] == boolean_false_node)
+	      fn_level = ix + 1;
+	    else if (fn_level < 0)
+	      fn_level = ix;
 	  }
+	
+	tree val = TREE_VALUE (pos);
+
+	dims[ix] = val ? TREE_INT_CST_LOW (val) : -1;
+	pos = TREE_CHAIN (pos);
       }
 
-    bool changed = targetm.goacc.validate_dims (current_function_decl, dims);
+    bool changed = targetm.goacc.validate_dims (current_function_decl,
+						dims, fn_level);
 
-    /* Default anything left undefaulted to 1.  */
+    /* Default anything left to 1.  */
     for (ix = 0; ix != GOMP_DIM_MAX; ix++)
       if (dims[ix] < 0)
 	{
-	  dims[ix] = (int)(dims[ix] < -1);
+	  dims[ix] = 1;
 	  changed = true;
 	}
-  
+
     if (changed)
       {
 	/* Replace the attribute with new values.  */
 	pos = NULL_TREE;
 	for (ix = GOMP_DIM_MAX; ix--;)
-	  pos = tree_cons (NULL_TREE,
+	  pos = tree_cons (purpose[ix],
 			   build_int_cst (integer_type_node, dims[ix]),
 			   pos);
 	replace_oacc_fn_attrib (current_function_decl, pos);
@@ -14722,7 +14777,8 @@  execute_oacc_transform ()
    hook.  */
 
 bool
-default_goacc_validate_dims (tree ARG_UNUSED (decl), int *ARG_UNUSED (dims))
+default_goacc_validate_dims (tree ARG_UNUSED (decl), int *ARG_UNUSED (dims),
+			     int ARG_UNUSED (fn_level))
 {
   bool changed = false;
 
Index: gcc/omp-low.h
===================================================================
--- gcc/omp-low.h	(revision 226860)
+++ gcc/omp-low.h	(working copy)
@@ -30,6 +30,8 @@  extern bool make_gimple_omp_edges (basic
 extern void omp_finish_file (void);
 extern bool gimple_stmt_omp_data_i_init_p (gimple);
 extern basic_block loop_get_oacc_kernels_region_entry (struct loop *);
+extern void replace_oacc_fn_attrib (tree, tree);
+extern tree build_oacc_routine_dims (tree);
 extern tree get_oacc_fn_attrib (tree);
 
 extern GTY(()) vec<tree, va_gc> *offload_funcs;
Index: gcc/target.def
===================================================================
--- gcc/target.def	(revision 226860)
+++ gcc/target.def	(working copy)
@@ -1651,7 +1651,7 @@  in anything that needs default to non-un
 Defaults are represented as -1.  Diagnostics should be issuedas \n\
 ppropriate.  Return true if changes have been made.  You must override\n\
 this hook to provide dimensions larger than 1.",
-bool, (tree, int []),
+bool, (tree, int [], int),
 default_goacc_validate_dims)
 
 DEFHOOK
Index: gcc/targhooks.h
===================================================================
--- gcc/targhooks.h	(revision 226860)
+++ gcc/targhooks.h	(working copy)
@@ -107,7 +107,7 @@  extern unsigned default_add_stmt_cost (v
 extern void default_finish_cost (void *, unsigned *, unsigned *, unsigned *);
 extern void default_destroy_cost_data (void *);
 
-extern bool default_goacc_validate_dims (tree, int []);
+extern bool default_goacc_validate_dims (tree, int [], int);
 extern unsigned default_goacc_dim_limit (unsigned);
 extern bool default_goacc_fork_join (gimple_stmt_iterator *, gimple,
 				     const int [], bool);
Index: gcc/testsuite/c-c++-common/goacc/routine-4.c
===================================================================
--- gcc/testsuite/c-c++-common/goacc/routine-4.c	(revision 226860)
+++ gcc/testsuite/c-c++-common/goacc/routine-4.c	(working copy)
@@ -1,87 +1,74 @@ 
 /* Test invalid use of clauses with routine.  */
 /* { dg-do compile } */
 
-#pragma acc routine gang worker /* { dg-error "invalid combination" } */
+#pragma acc routine gang worker /* { dg-error "multiple loop axes" } */
 void
 f1 (void)
 {
 }
 
-#pragma acc routine worker gang /* { dg-error "invalid combination" } */
+#pragma acc routine worker gang /* { dg-error "multiple loop axes" } */
 void
 f1a (void)
 {
 }
 
-#pragma acc routine gang vector /* { dg-error "invalid combination" } */
+#pragma acc routine gang vector /* { dg-error "multiple loop axes" } */
 void
 f2 (void)
 {
 }
 
-#pragma acc routine vector gang /* { dg-error "invalid combination" } */
+#pragma acc routine vector gang /* { dg-error "multiple loop axes" } */
 void
 f2a (void)
 {
 }
 
-#pragma acc routine gang seq /* { dg-error "invalid combination" } */
+#pragma acc routine gang seq /* { dg-error "multiple loop axes" } */
 void
 f3 (void)
 {
 }
 
-#pragma acc routine seq gang /* { dg-error "invalid combination" } */
+#pragma acc routine seq gang /* { dg-error "multiple loop axes" } */
 void
 f3a (void)
 {
 }
 
-#pragma acc routine worker vector /* { dg-error "invalid combination" } */
+#pragma acc routine worker vector /* { dg-error "multiple loop axes" } */
 void
 f4 (void)
 {
 }
 
-#pragma acc routine vector worker /* { dg-error "invalid combination" } */
+#pragma acc routine vector worker /* { dg-error "multiple loop axes" } */
 void
 f4a (void)
 {
 }
 
-#pragma acc routine worker seq /* { dg-error "invalid combination" } */
+#pragma acc routine worker seq /* { dg-error "multiple loop axes" } */
 void
 f5 (void)
 {
 }
 
-#pragma acc routine seq worker /* { dg-error "invalid combination" } */
+#pragma acc routine seq worker /* { dg-error "multiple loop axes" } */
 void
 f5a (void)
 {
 }
 
-#pragma acc routine vector seq /* { dg-error "invalid combination" } */
+#pragma acc routine vector seq /* { dg-error "multiple loop axes" } */
 void
 f6 (void)
 {
 }
 
-#pragma acc routine seq vector /* { dg-error "invalid combination" } */
+#pragma acc routine seq vector /* { dg-error "multiple loop axes" } */
 void
 f6a (void)
 {
 }
-
-#pragma acc routine (g1) gang worker /* { dg-error "invalid combination" } */
-#pragma acc routine (g2) worker gang /* { dg-error "invalid combination" } */
-#pragma acc routine (g3) gang vector /* { dg-error "invalid combination" } */
-#pragma acc routine (g4) vector gang /* { dg-error "invalid combination" } */
-#pragma acc routine (g5) gang seq /* { dg-error "invalid combination" } */
-#pragma acc routine (g6) seq gang /* { dg-error "invalid combination" } */
-#pragma acc routine (g7) worker vector /* { dg-error "invalid combination" } */
-#pragma acc routine (g8) vector worker /* { dg-error "invalid combination" } */
-#pragma acc routine (g9) worker seq /* { dg-error "invalid combination" } */
-#pragma acc routine (g10) seq worker /* { dg-error "invalid combination" } */
-#pragma acc routine (g11) vector seq /* { dg-error "invalid combination" } */
-#pragma acc routine (g12) seq vector /* { dg-error "invalid combination" } */