2016-11-22  Cesar Philippidis  <cesar@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

	gcc/c-family/
	* c-attribs.c (c_common_attribute_table): Adjust "omp declare target".
	* c-pragma.h (enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_BIND
	and PRAGMA_OACC_CLAUSE_NOHOST.

	gcc/
	* gimplify.c (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_BIND and
	OMP_CLAUSE_NOHOST.
	(gimplify_adjust_omp_clauses): Likewise.
	* omp-low.c (scan_sharing_clauses): Likewise.
	(verify_oacc_routine_clauses): New function.
	(maybe_discard_oacc_function): New function.
	(execute_oacc_device_lower): Don't generate code for NOHOST.
	* omp-low.h (verify_oacc_routine_clauses): Declare.
	* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_BIND and
	OMP_CLAUSE_NOHOST.
	* tree-pretty-print.c (dump_omp_clause): Likewise.
	* tree.c (omp_clause_num_ops): Likewise.
	(omp_clause_code_name): Likewise.
	(walk_tree_1): Handle OMP_CLAUSE_BIND, OMP_CLAUSE_NOHOST.
	* tree.h (OMP_CLAUSE_BIND_NAME): Define.


diff --git a/gcc/c-family/c-attribs.c b/gcc/c-family/c-attribs.c
index 964efe9..49999b8 100644
--- a/gcc/c-family/c-attribs.c
+++ b/gcc/c-family/c-attribs.c
@@ -322,7 +322,7 @@ const struct attribute_spec c_common_attribute_table[] =
 			      handle_omp_declare_simd_attribute, false },
   { "simd",		      0, 1, true,  false, false,
 			      handle_simd_attribute, false },
-  { "omp declare target",     0, 0, true, false, false,
+  { "omp declare target",     0, -1, true, false, false,
 			      handle_omp_declare_target_attribute, false },
   { "omp declare target link", 0, 0, true, false, false,
 			      handle_omp_declare_target_attribute, false },
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index 6d9cb08..dd2722a 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -149,6 +149,7 @@ enum pragma_omp_clause {
   /* Clauses for OpenACC.  */
   PRAGMA_OACC_CLAUSE_ASYNC = PRAGMA_CILK_CLAUSE_VECTORLENGTH + 1,
   PRAGMA_OACC_CLAUSE_AUTO,
+  PRAGMA_OACC_CLAUSE_BIND,
   PRAGMA_OACC_CLAUSE_COPY,
   PRAGMA_OACC_CLAUSE_COPYOUT,
   PRAGMA_OACC_CLAUSE_CREATE,
@@ -158,6 +159,7 @@ enum pragma_omp_clause {
   PRAGMA_OACC_CLAUSE_GANG,
   PRAGMA_OACC_CLAUSE_HOST,
   PRAGMA_OACC_CLAUSE_INDEPENDENT,
+  PRAGMA_OACC_CLAUSE_NOHOST,
   PRAGMA_OACC_CLAUSE_NUM_GANGS,
   PRAGMA_OACC_CLAUSE_NUM_WORKERS,
   PRAGMA_OACC_CLAUSE_PRESENT,
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 8611060..04b591e 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8373,6 +8373,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 ();
 	}
@@ -9112,6 +9114,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	  remove = true;
 	  break;
 
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	default:
 	  gcc_unreachable ();
 	}
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 7c58c03..b8a414b 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2201,6 +2201,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	    install_var_local (decl, ctx);
 	  break;
 
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE__CACHE_:
 	default:
@@ -2365,6 +2367,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
 	case OMP_CLAUSE__GRIDDIM_:
 	  break;
 
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	case OMP_CLAUSE_TILE:
 	case OMP_CLAUSE__CACHE_:
 	default:
@@ -12684,9 +12688,192 @@ set_oacc_fn_attrib (tree fn, tree clauses, bool is_kernel, vec<tree> *args)
     }
 }
 
-/*  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
+/* Verify OpenACC routine clauses.
+
+   Returns 0 if FNDECL should be marked as an accelerator routine, 1 if it has
+   already been marked in compatible way, and -1 if incompatible.  Upon
+   returning, the chain of clauses will contain exactly one clause specifying
+   the level of parallelism.  */
+
+int
+verify_oacc_routine_clauses (tree fndecl, tree *clauses, location_t loc,
+			     const char *routine_str)
+{
+  tree c_level = NULL_TREE;
+  tree c_bind = NULL_TREE;
+  tree c_nohost = NULL_TREE;
+  tree c_p = NULL_TREE;
+  for (tree c = *clauses; c; c_p = 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:
+	if (c_level == NULL_TREE)
+	  c_level = c;
+	else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_CODE (c_level))
+	  {
+	    /* This has already been diagnosed in the front ends.  */
+	    /* Drop the duplicate clause.  */
+	    gcc_checking_assert (c_p != NULL_TREE);
+	    OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
+	    c = c_p;
+	  }
+	else
+	  {
+	    error_at (OMP_CLAUSE_LOCATION (c),
+		      "%qs specifies a conflicting level of parallelism",
+		      omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+	    inform (OMP_CLAUSE_LOCATION (c_level),
+		    "previous %qs clause",
+		    omp_clause_code_name[OMP_CLAUSE_CODE (c_level)]);
+	    /* Drop the conflicting clause.  */
+	    gcc_checking_assert (c_p != NULL_TREE);
+	    OMP_CLAUSE_CHAIN (c_p) = OMP_CLAUSE_CHAIN (c);
+	    c = c_p;
+	  }
+	break;
+      case OMP_CLAUSE_BIND:
+	/* Don't bother with duplicate clauses at this point.  */
+	c_bind = c;
+	break;
+      case OMP_CLAUSE_NOHOST:
+	/* Don't bother with duplicate clauses at this point.  */
+	c_nohost = c;
+	break;
+      default:
+	gcc_unreachable ();
+      }
+  if (c_level == NULL_TREE)
+    {
+      /* OpenACC 2.5 makes this an error; for the current OpenACC 2.0a
+	 implementation add an implicit "seq" clause.  */
+      c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
+      OMP_CLAUSE_CHAIN (c_level) = *clauses;
+      *clauses = c_level;
+    }
+  /* In *clauses, we now have exactly one clause specifying the level of
+     parallelism.  */
+
+  /* Still got some work to do for Fortran...  */
+  if (fndecl == NULL_TREE)
+    return 0;
+
+  tree attr
+    = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl));
+  if (attr != NULL_TREE)
+    {
+      /* If a "#pragma acc routine" has already been applied, just verify
+	 this one for compatibility.  */
+      /* Collect previous directive's clauses.  */
+      tree c_level_p = NULL_TREE;
+      tree c_bind_p = NULL_TREE;
+      tree c_nohost_p = NULL_TREE;
+      for (tree c = TREE_VALUE (attr); 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:
+	    gcc_checking_assert (c_level_p == NULL_TREE);
+	    c_level_p = c;
+	    break;
+	  case OMP_CLAUSE_BIND:
+	    /* Don't bother with duplicate clauses at this point.  */
+	    c_bind_p = c;
+	    break;
+	  case OMP_CLAUSE_NOHOST:
+	    /* Don't bother with duplicate clauses at this point.  */
+	    c_nohost_p = c;
+	    break;
+	  default:
+	    gcc_unreachable ();
+	  }
+      gcc_checking_assert (c_level_p != NULL_TREE);
+      /* ..., and compare to current directive's, which we've already collected
+	 above.  */
+      tree c_diag;
+      tree c_diag_p;
+      /* Matching level of parallelism?  */
+      if (OMP_CLAUSE_CODE (c_level) != OMP_CLAUSE_CODE (c_level_p))
+	{
+	  c_diag = c_level;
+	  c_diag_p = c_level_p;
+	  goto incompatible;
+	}
+      /* Matching bind clauses?  */
+      if ((c_bind == NULL_TREE) != (c_bind_p == NULL_TREE))
+	{
+	  c_diag = c_bind;
+	  c_diag_p = c_bind_p;
+	  goto incompatible;
+	}
+      /* Matching bind clauses' names?  */
+      if ((c_bind != NULL_TREE) && (c_bind_p != NULL_TREE))
+	{
+	  tree c_bind_name = OMP_CLAUSE_BIND_NAME (c_bind);
+	  tree c_bind_name_p = OMP_CLAUSE_BIND_NAME (c_bind_p);
+	  /* TODO: will/should actually be the trees/strings/string pointers be
+	     identical?  */
+	  if (strcmp (TREE_STRING_POINTER (c_bind_name),
+		      TREE_STRING_POINTER (c_bind_name_p)) != 0)
+	    {
+	      c_diag = c_bind;
+	      c_diag_p = c_bind_p;
+	      goto incompatible;
+	    }
+	}
+      /* Matching nohost clauses?  */
+      if ((c_nohost == NULL_TREE) != (c_nohost_p == NULL_TREE))
+	{
+	  c_diag = c_nohost;
+	  c_diag_p = c_nohost_p;
+	  goto incompatible;
+	}
+      /* Compatible.  */
+      return 1;
+
+    incompatible:
+      if (c_diag != NULL_TREE)
+	error_at (OMP_CLAUSE_LOCATION (c_diag),
+		  "incompatible %qs clause when applying"
+		  " %<%s%> to %qD, which has already been"
+		  " marked as an accelerator routine",
+		  omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)],
+		  routine_str, fndecl);
+      else if (c_diag_p != NULL_TREE)
+	error_at (loc,
+		  "missing %qs clause when applying"
+		  " %<%s%> to %qD, which has already been"
+		  " marked as an accelerator routine",
+		  omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)],
+		  routine_str, fndecl);
+      else
+	gcc_unreachable ();
+      if (c_diag_p != NULL_TREE)
+	inform (OMP_CLAUSE_LOCATION (c_diag_p),
+		"previous %qs clause",
+		omp_clause_code_name[OMP_CLAUSE_CODE (c_diag_p)]);
+      else
+	{
+	  /* In the front ends, we don't preserve location information for the
+	     OpenACC routine directive itself.  However, that of c_level_p
+	     should be close.  */
+	  location_t loc_routine = OMP_CLAUSE_LOCATION (c_level_p);
+	  inform (loc_routine, "previous %qs clause near to here",
+		  omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]);
+	}
+      /* Incompatible.  */
+      return -1;
+    }
+
+  return 0;
+}
+
+/*  Process the OpenACC routine's clauses to generate an attribute
+    for the level of parallelism.  All dimensions have a size of zero
     (dynamic).  TREE_PURPOSE is set to indicate whether that dimension
     can have a loop partitioned on it.  non-zero indicates
     yes, zero indicates no.  By construction once a non-zero has been
@@ -19694,6 +19881,28 @@ default_goacc_reduction (gcall *call)
   gsi_replace_with_seq (&gsi, seq, true);
 }
 
+/* Determine whether DECL should be discarded in this offload
+   compilation.  */
+
+static bool
+maybe_discard_oacc_function (tree decl)
+{
+  tree attr = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl));
+
+  if (!attr)
+    return false;
+
+  enum omp_clause_code kind = OMP_CLAUSE_NOHOST;
+  
+#ifdef ACCEL_COMPILER
+  kind = OMP_CLAUSE_BIND;
+#endif
+  if (find_omp_clause (TREE_VALUE (attr), kind))
+    return true;
+
+  return false;
+}
+
 /* Main entry point for oacc transformations which run on the device
    compiler after LTO, so we know what the target device is at this
    point (including the host fallback).  */
@@ -19707,6 +19916,14 @@ execute_oacc_device_lower ()
     /* Not an offloaded function.  */
     return 0;
 
+  if (maybe_discard_oacc_function (current_function_decl))
+    {
+      if (dump_file)
+	fprintf (dump_file, "Discarding function\n");
+      TREE_ASM_WRITTEN (current_function_decl) = 1;
+      return TODO_discard_function;
+    }
+
   /* Parse the default dim argument exactly once.  */
   if ((const void *)flag_openacc_dims != &flag_openacc_dims)
     {
diff --git a/gcc/omp-low.h b/gcc/omp-low.h
index b1f7885..2602a12 100644
--- a/gcc/omp-low.h
+++ b/gcc/omp-low.h
@@ -31,6 +31,7 @@ extern bool make_gimple_omp_edges (basic_block, struct omp_region **, int *);
 extern void omp_finish_file (void);
 extern tree omp_member_access_dummy_var (tree);
 extern void replace_oacc_fn_attrib (tree, tree);
+extern int verify_oacc_routine_clauses (tree, tree *, location_t, const char *);
 extern tree build_oacc_routine_dims (tree);
 extern tree get_oacc_fn_attrib (tree);
 extern void set_oacc_fn_attrib (tree, tree, bool, vec<tree> *);
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index a3d220d..bd0e254 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -465,7 +465,13 @@ enum omp_clause_code {
 
   /* OpenMP internal-only clause to specify grid dimensions of a gridified
      kernel.  */
-  OMP_CLAUSE__GRIDDIM_
+  OMP_CLAUSE__GRIDDIM_,
+
+  /* OpenACC clause: bind (string).  */
+  OMP_CLAUSE_BIND,
+
+  /* OpenACC clause: nohost.  */
+  OMP_CLAUSE_NOHOST
 };
 
 #undef DEFTREESTRUCT
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 096eefd..5494441 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -1022,6 +1022,15 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
 			 spc, flags, false);
       pp_right_paren (pp);
       break;
+    case OMP_CLAUSE_NOHOST:
+      pp_string (pp, "nohost");
+      break;
+    case OMP_CLAUSE_BIND:
+      pp_string (pp, "bind(");
+      dump_generic_node (pp, OMP_CLAUSE_BIND_NAME (clause),
+			 spc, flags, false);
+      pp_string (pp, ")");
+      break;
 
     case OMP_CLAUSE__GRIDDIM_:
       pp_string (pp, "_griddim_(");
diff --git a/gcc/tree.c b/gcc/tree.c
index a4c5b1b..521fb22 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -329,6 +329,8 @@ unsigned const char omp_clause_num_ops[] =
   1, /* OMP_CLAUSE_VECTOR_LENGTH  */
   1, /* OMP_CLAUSE_TILE  */
   2, /* OMP_CLAUSE__GRIDDIM_  */
+  1, /* OMP_CLAUSE_BIND  */
+  0, /* OMP_CLAUSE_NOHOST  */
 };
 
 const char * const omp_clause_code_name[] =
@@ -399,7 +401,9 @@ const char * const omp_clause_code_name[] =
   "num_workers",
   "vector_length",
   "tile",
-  "_griddim_"
+  "_griddim_",
+  "bind",
+  "nohost",
 };
 
 
@@ -11871,6 +11875,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
 	case OMP_CLAUSE__LOOPTEMP_:
 	case OMP_CLAUSE__SIMDUID_:
 	case OMP_CLAUSE__CILK_FOR_COUNT_:
+	case OMP_CLAUSE_BIND:
 	  WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 0));
 	  /* FALLTHRU */
 
@@ -11892,6 +11897,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
 	case OMP_CLAUSE_DEFAULTMAP:
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_NOHOST:
 	case OMP_CLAUSE_TILE:
 	  WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
 
diff --git a/gcc/tree.h b/gcc/tree.h
index b4ec3fd..a17606f 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1526,6 +1526,9 @@ extern void protected_set_expr_location (tree, location_t);
 #define OMP_CLAUSE_VECTOR_LENGTH_EXPR(NODE) \
   OMP_CLAUSE_OPERAND ( \
     OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_VECTOR_LENGTH), 0)
+#define OMP_CLAUSE_BIND_NAME(NODE) \
+  OMP_CLAUSE_OPERAND ( \
+    OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_BIND), 0)
 
 #define OMP_CLAUSE_DEPEND_KIND(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEPEND)->omp_clause.subcode.depend_kind)
