diff mbox

[4/5] OpenACC tile clause support, Fortran front-end parts

Message ID cbe5aca8-77c5-0057-1381-4a99e43b4318@codesourcery.com
State New
Headers show

Commit Message

Cesar Philippidis Nov. 12, 2016, 4:51 p.m. UTC
On 11/11/2016 02:34 AM, Jakub Jelinek wrote:
> On Thu, Nov 10, 2016 at 06:46:46PM +0800, Chung-Lin Tang wrote:


And here's the patch.

Cesar

>> 2016-XX-XX  Cesar Philippidis  <cesar@codesourcery.com>

>>

>> 	fortran/

>> 	* openmp.c (resolve_oacc_positive_int_expr): Promote the warning

>> 	to an error.

>> 	(resolve_oacc_loop_blocks): Use integer zero to represent the '*'

>> 	tile argument.

>> 	(resolve_omp_clauses): Error on directives containing both tile

>> 	and collapse clauses.

>> 	* trans-openmp.c (gfc_trans_omp_do): Lower tiled loops like

>> 	collapsed loops.

>>

>>

>> 	gcc/testsuite/

>> 	* gfortran.dg/goacc/loop-2.f95: Change expected tile clause

>>         warnings to errors.

>>         * gfortran.dg/goacc/loop-5.f95: Likewise.

>>         * gfortran.dg/goacc/sie.f95: Likewise.

>>         * gfortran.dg/goacc/tile-1.f90: New test.

>>         * gfortran.dg/goacc/tile-2.f90: New test

>>         * gfortran.dg/goacc/tile-lowering.f95: New test.

> 

> Again, 8 spaces in ChangeLog.  Missing full stop after New test

> 

>> --- fortran/openmp.c	(revision 241809)

>> +++ fortran/openmp.c	(working copy)

>> @@ -3024,8 +3024,8 @@ resolve_oacc_positive_int_expr (gfc_expr *expr, co

>>    resolve_oacc_scalar_int_expr (expr, clause);

>>    if (expr->expr_type == EXPR_CONSTANT && expr->ts.type == BT_INTEGER

>>        && mpz_sgn(expr->value.integer) <= 0)

>> -    gfc_warning (0, "INTEGER expression of %s clause at %L must be positive",

>> -		     clause, &expr->where);

>> +    gfc_error ("INTEGER expression of %s clause at %L must be positive",

>> +	       clause, &expr->where);

>>  }

> 

> This can't be against current trunk.  The current routine is shared with

> OpenMP and gfc_error is undesirable there.

> 

>> @@ -3859,6 +3859,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus

>>      if (omp_clauses->wait_list)

>>        for (el = omp_clauses->wait_list; el; el = el->next)

>>  	resolve_oacc_scalar_int_expr (el->expr, "WAIT");

>> +  if (omp_clauses->collapse && omp_clauses->tile_list)

>> +    gfc_error ("Incompatible use of TILE and COLLAPSE at %L", &code->loc);

> 

> Shouldn't you in that case for error recovery clear collapse (or tile_list)?

> 

> 	Jakub

>

Comments

Jakub Jelinek Nov. 18, 2016, 11:24 a.m. UTC | #1
On Sat, Nov 12, 2016 at 08:51:00AM -0800, Cesar Philippidis wrote:
> On 11/11/2016 02:34 AM, Jakub Jelinek wrote:

> > On Thu, Nov 10, 2016 at 06:46:46PM +0800, Chung-Lin Tang wrote:

> 

> And here's the patch.


The patch doesn't look like OpenACC tile clause fortran support,
but bind/nohost clause C/C++ support.

	Jakub
diff mbox

Patch

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

	gcc/cp/
	* cp-tree.h (bind_decls_match): Declare.
	* decl.c (bind_decls_match): New function.
	* parser.c (cp_parser_omp_clause_name): 
	(cp_parser_oacc_shape_clause): New location_t loc argument.  Use it
	to report more accurate diagnostics.
	(cp_parser_oacc_clause_bind): New function.
	(cp_parser_oacc_all_clauses): Handle OpenACC bind and nohost clauses.
	Update calls to c_parser_oacc_{simple,shape}_clause.
	(OACC_ROUTINE_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{BIND,NOHOST}.
	(cp_parser_oacc_routine): Update diagnostics.
	(cp_parser_late_parsing_oacc_routine): Likewise.
	(cp_finalize_oacc_routine): Likewise.
	* semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_{BIND,NOHOST}.


diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 8183775..efd5450 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -5772,6 +5772,7 @@  extern void finish_scope			(void);
 extern void push_switch				(tree);
 extern void pop_switch				(void);
 extern tree make_lambda_name			(void);
+extern int bind_decls_match			(tree, tree);
 extern int decls_match				(tree, tree);
 extern tree duplicate_decls			(tree, tree, bool);
 extern tree declare_local_label			(tree);
diff --git a/gcc/cp/decl.c b/gcc/cp/decl.c
index 185c98b..7b242781 100644
--- a/gcc/cp/decl.c
+++ b/gcc/cp/decl.c
@@ -1198,6 +1198,138 @@  decls_match (tree newdecl, tree olddecl)
   return types_match;
 }
 
+/* Similiar to decls_match, but only applies to FUNCTION_DECLS.  Functions
+   in separate namespaces may match.
+*/
+
+int
+bind_decls_match (tree newdecl, tree olddecl)
+{
+  int types_match;
+
+  if (newdecl == olddecl)
+    return 1;
+
+  if (TREE_CODE (newdecl) != TREE_CODE (olddecl))
+    /* If the two DECLs are not even the same kind of thing, we're not
+       interested in their types.  */
+    return 0;
+
+  gcc_assert (DECL_P (newdecl));
+  gcc_assert (TREE_CODE (newdecl) == FUNCTION_DECL);
+
+  tree f1 = TREE_TYPE (newdecl);
+  tree f2 = TREE_TYPE (olddecl);
+  tree p1 = TYPE_ARG_TYPES (f1);
+  tree p2 = TYPE_ARG_TYPES (f2);
+  tree r2;
+
+  /* Specializations of different templates are different functions
+     even if they have the same type.  */
+  tree t1 = (DECL_USE_TEMPLATE (newdecl)
+	     ? DECL_TI_TEMPLATE (newdecl)
+	     : NULL_TREE);
+  tree t2 = (DECL_USE_TEMPLATE (olddecl)
+	     ? DECL_TI_TEMPLATE (olddecl)
+	     : NULL_TREE);
+  if (t1 != t2)
+    return 0;
+
+  if (CP_DECL_CONTEXT (newdecl) != CP_DECL_CONTEXT (olddecl)
+      && TREE_CODE (CP_DECL_CONTEXT (newdecl)) != NAMESPACE_DECL
+      && TREE_CODE (CP_DECL_CONTEXT (olddecl)) != NAMESPACE_DECL
+      && ! (DECL_EXTERN_C_P (newdecl)
+	    && DECL_EXTERN_C_P (olddecl)))
+    return 0;
+
+  /* A new declaration doesn't match a built-in one unless it
+     is also extern "C".  */
+  if (DECL_IS_BUILTIN (olddecl)
+      && DECL_EXTERN_C_P (olddecl) && !DECL_EXTERN_C_P (newdecl))
+    return 0;
+
+  if (TREE_CODE (f1) != TREE_CODE (f2))
+    return 0;
+
+  /* A declaration with deduced return type should use its pre-deduction
+     type for declaration matching.  */
+  r2 = fndecl_declared_return_type (olddecl);
+
+  if (same_type_p (TREE_TYPE (f1), r2))
+    {
+      if (!prototype_p (f2) && DECL_EXTERN_C_P (olddecl)
+	  && (DECL_BUILT_IN (olddecl)
+#ifndef NO_IMPLICIT_EXTERN_C
+	      || (DECL_IN_SYSTEM_HEADER (newdecl) && !DECL_CLASS_SCOPE_P (newdecl))
+	      || (DECL_IN_SYSTEM_HEADER (olddecl) && !DECL_CLASS_SCOPE_P (olddecl))
+#endif
+	      ))
+	{
+	  types_match = self_promoting_args_p (p1);
+	  if (p1 == void_list_node)
+	    TREE_TYPE (newdecl) = TREE_TYPE (olddecl);
+	}
+#ifndef NO_IMPLICIT_EXTERN_C
+      else if (!prototype_p (f1)
+	       && (DECL_EXTERN_C_P (olddecl)
+		   && DECL_IN_SYSTEM_HEADER (olddecl)
+		   && !DECL_CLASS_SCOPE_P (olddecl))
+	       && (DECL_EXTERN_C_P (newdecl)
+		   && DECL_IN_SYSTEM_HEADER (newdecl)
+		   && !DECL_CLASS_SCOPE_P (newdecl)))
+	{
+	  types_match = self_promoting_args_p (p2);
+	  TREE_TYPE (newdecl) = TREE_TYPE (olddecl);
+	}
+#endif
+      else
+	types_match =
+	  compparms (p1, p2)
+	  && type_memfn_rqual (f1) == type_memfn_rqual (f2)
+	  && (TYPE_ATTRIBUTES (TREE_TYPE (newdecl)) == NULL_TREE
+	      || comp_type_attributes (TREE_TYPE (newdecl),
+				       TREE_TYPE (olddecl)) != 0);
+    }
+  else
+    types_match = 0;
+
+  /* The decls dont match if they correspond to two different versions
+     of the same function.   Disallow extern "C" functions to be
+     versions for now.  */
+  if (types_match
+      && !DECL_EXTERN_C_P (newdecl)
+      && !DECL_EXTERN_C_P (olddecl)
+      && targetm.target_option.function_versions (newdecl, olddecl))
+    {
+      /* Mark functions as versions if necessary.  Modify the mangled decl
+	 name if necessary.  */
+      if (DECL_FUNCTION_VERSIONED (newdecl)
+	  && DECL_FUNCTION_VERSIONED (olddecl))
+	return 0;
+      if (!DECL_FUNCTION_VERSIONED (newdecl))
+	{
+	  DECL_FUNCTION_VERSIONED (newdecl) = 1;
+	  if (DECL_ASSEMBLER_NAME_SET_P (newdecl))
+	    mangle_decl (newdecl);
+	}
+      if (!DECL_FUNCTION_VERSIONED (olddecl))
+	{
+	  DECL_FUNCTION_VERSIONED (olddecl) = 1;
+	  if (DECL_ASSEMBLER_NAME_SET_P (olddecl))
+	    mangle_decl (olddecl);
+	}
+      cgraph_node::record_function_versions (olddecl, newdecl);
+      return 0;
+    }
+
+  /* Normal functions can be constrained, as can variable partial
+     specializations.  */
+  if (types_match && VAR_OR_FUNCTION_DECL_P (newdecl))
+    types_match = equivalently_constrained (newdecl, olddecl);
+
+  return types_match;
+}
+
 /* If NEWDECL is `static' and an `extern' was seen previously,
    warn about it.  OLDDECL is the previous declaration.
 
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 7b95dba..33ee6a5 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -30223,6 +30223,10 @@  cp_parser_omp_clause_name (cp_parser *parser)
 	  else if (!strcmp ("async", p))
 	    result = PRAGMA_OACC_CLAUSE_ASYNC;
 	  break;
+	case 'b':
+	  if (!strcmp ("bind", p))
+	    result = PRAGMA_OACC_CLAUSE_BIND;
+	  break;
 	case 'c':
 	  if (!strcmp ("collapse", p))
 	    result = PRAGMA_OMP_CLAUSE_COLLAPSE;
@@ -30302,6 +30306,8 @@  cp_parser_omp_clause_name (cp_parser *parser)
 	    result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
 	  else if (!strcmp ("nowait", p))
 	    result = PRAGMA_OMP_CLAUSE_NOWAIT;
+	  else if (!strcmp ("nohost", p))
+	    result = PRAGMA_OACC_CLAUSE_NOHOST;
 	  else if (flag_cilkplus && !strcmp ("nomask", p))
 	    result = PRAGMA_CILK_CLAUSE_NOMASK;
 	  else if (!strcmp ("num_gangs", p))
@@ -30778,13 +30784,13 @@  cp_parser_oacc_single_int_clause (cp_parser *parser, omp_clause_code code,
 */
 
 static tree
-cp_parser_oacc_shape_clause (cp_parser *parser, omp_clause_code kind,
+cp_parser_oacc_shape_clause (cp_parser *parser, location_t loc,
+			     omp_clause_code kind,
 			     const char *str, tree list)
 {
   const char *id = "num";
   cp_lexer *lexer = parser->lexer;
   tree ops[2] = { NULL_TREE, NULL_TREE }, c;
-  location_t loc = cp_lexer_peek_token (lexer)->location;
 
   if (kind == OMP_CLAUSE_VECTOR)
     id = "length";
@@ -32571,6 +32577,82 @@  cp_parser_oacc_clause_async (cp_parser *parser, tree list)
   return list;
 }
 
+/* OpenACC 2.0:
+   bind ( identifier )
+   bind ( string-literal ) */
+
+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;
+    }
+  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 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 (!decl || decl == error_mark_node)
+	error_at (token->location, "%qE has not been declared",
+		  token->u.value);
+      else if (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 (TREE_CODE (decl) != FUNCTION_DECL)
+	error_at (token->location,
+		  "%qE does not refer to a function",
+		  token->u.value);
+      else
+	name = decl;
+    }
+  else if (cp_lexer_next_token_is (parser->lexer, CPP_STRING))
+    {
+      name = token->u.value;
+      cp_lexer_consume_token (parser->lexer);
+
+      /* This shouldn't be an empty string.  */
+      if (strcmp (TREE_STRING_POINTER (name), "\"\"") == 0)
+	error_at (token->location,
+		  "bind argument must not be an empty string");
+
+      parser->translate_strings_p = save_translate_strings_p;
+    }
+  else
+    {
+      cp_parser_error (parser,
+		       "expected identifier or character string literal");
+      cp_parser_skip_to_closing_parenthesis (parser, false, false, true);
+      return list;
+    }
+  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;
+}
+
 /* Parse all OpenACC clauses.  The set clauses allowed by the directive
    is a bitmask in MASK.  Return the list of clauses found.  */
 
@@ -32607,6 +32689,10 @@  cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 						 clauses, here);
 	  c_name = "auto";
 	  break;
+	case PRAGMA_OACC_CLAUSE_BIND:
+	  clauses = cp_parser_oacc_clause_bind (parser, clauses);
+	  c_name = "bind";
+	  break;
 	case PRAGMA_OACC_CLAUSE_COLLAPSE:
 	  clauses = cp_parser_omp_clause_collapse (parser, clauses, here);
 	  c_name = "collapse";
@@ -32654,7 +32740,7 @@  cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  break;
 	case PRAGMA_OACC_CLAUSE_GANG:
 	  c_name = "gang";
-	  clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG,
+	  clauses = cp_parser_oacc_shape_clause (parser, here, OMP_CLAUSE_GANG,
 						 c_name, clauses);
 	  break;
 	case PRAGMA_OACC_CLAUSE_HOST:
@@ -32675,6 +32761,11 @@  cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
 	  c_name = "link";
 	  break;
+	case PRAGMA_OACC_CLAUSE_NOHOST:
+	  clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_NOHOST,
+						  clauses, here);
+	  c_name = "nohost";
+	  break;
 	case PRAGMA_OACC_CLAUSE_NUM_GANGS:
 	  code = OMP_CLAUSE_NUM_GANGS;
 	  c_name = "num_gangs";
@@ -32736,7 +32827,8 @@  cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  break;
 	case PRAGMA_OACC_CLAUSE_VECTOR:
 	  c_name = "vector";
-	  clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
+	  clauses = cp_parser_oacc_shape_clause (parser, here,
+						 OMP_CLAUSE_VECTOR,
 						 c_name, clauses);
 	  break;
 	case PRAGMA_OACC_CLAUSE_VECTOR_LENGTH:
@@ -32751,7 +32843,8 @@  cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
 	  break;
 	case PRAGMA_OACC_CLAUSE_WORKER:
 	  c_name = "worker";
-	  clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_WORKER,
+	  clauses = cp_parser_oacc_shape_clause (parser, here,
+						 OMP_CLAUSE_WORKER,
 						 c_name, clauses);
 	  break;
 	default:
@@ -36992,7 +37085,10 @@  cp_parser_omp_taskloop (cp_parser *parser, cp_token *pragma_tok,
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR)		\
-	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ))
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_BIND)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST))
+
 
 
 /* Parse the OpenACC routine pragma.  This has an optional '( name )'
@@ -37051,6 +37147,9 @@  cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
 	= cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
 				      "#pragma acc routine",
 				      cp_lexer_peek_token (parser->lexer));
+      /* The clauses are in reverse order; fix that to make later diagnostic
+	 emission easier.  */
+      data.clauses = nreverse (data.clauses);
 
       if (decl && is_overloaded_fn (decl)
 	  && (TREE_CODE (decl) != FUNCTION_DECL
@@ -37062,16 +37161,6 @@  cp_parser_oacc_routine (cp_parser *parser, cp_token *pragma_tok,
 	  return;
 	}
 
-      /* Perhaps we should use the same rule as declarations in different
-	 namespaces?  */
-      if (!DECL_NAMESPACE_SCOPE_P (decl))
-	{
-	  error_at (name_loc,
-		    "%qD does not refer to a namespace scope function", decl);
-	  parser->oacc_routine = NULL;
-	  return;
-	}
-
       if (TREE_CODE (decl) != FUNCTION_DECL)
 	{
 	  error_at (name_loc, "%qD does not refer to a function", decl);
@@ -37147,6 +37236,9 @@  cp_parser_late_parsing_oacc_routine (cp_parser *parser, tree attrs)
   parser->oacc_routine->clauses
     = cp_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
 				  "#pragma acc routine", pragma_tok);
+  /* The clauses are in reverse order; fix that to make later diagnostic
+     emission easier.  */
+  parser->oacc_routine->clauses = nreverse (parser->oacc_routine->clauses);
   cp_parser_pop_lexer (parser);
   /* Later, cp_finalize_oacc_routine will process the clauses, and then set
      fndecl_seen.  */
@@ -37181,31 +37273,70 @@  cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
 	  return;
 	}
 
-      if (get_oacc_fn_attrib (fndecl))
+      /* Process the bind clause, if present.  */
+      for (tree c = parser->oacc_routine->clauses;
+	   c;
+	   c = OMP_CLAUSE_CHAIN (c))
 	{
-	  error_at (parser->oacc_routine->loc,
-		    "%<#pragma acc routine%> already applied to %qD", fndecl);
-	  parser->oacc_routine = NULL;
-	  return;
+	  if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_BIND)
+	    continue;
+
+	  tree bind_decl = OMP_CLAUSE_BIND_NAME (c);
+
+	  /* String arguments don't require any special treatment.  */
+	  if (TREE_CODE (bind_decl) != FUNCTION_DECL)
+	    break;
+
+	  if (!bind_decls_match (bind_decl, fndecl))
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"bind identifier %qE is not compatible with "
+			"function %qE", bind_decl, fndecl);
+	      parser->oacc_routine = NULL;
+	      return;
+	    }
+
+	  tree name_id = decl_assembler_name (bind_decl);
+	  tree name = build_string (IDENTIFIER_LENGTH (name_id),
+				    IDENTIFIER_POINTER (name_id));
+	  OMP_CLAUSE_BIND_NAME (c) = name;
+
+	  break;
 	}
 
-      if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+      int compatible
+	= verify_oacc_routine_clauses (fndecl, &parser->oacc_routine->clauses,
+				       parser->oacc_routine->loc,
+				       "#pragma acc routine");
+      if (compatible < 0)
 	{
-	  error_at (parser->oacc_routine->loc,
-		    "%<#pragma acc routine%> must be applied before %s",
-		    TREE_USED (fndecl) ? "use" : "definition");
 	  parser->oacc_routine = NULL;
 	  return;
 	}
+      if (compatible > 0)
+	{
+	}
+      else
+	{
+	  if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+	    {
+	      error_at (parser->oacc_routine->loc,
+			"%<#pragma acc routine%> must be applied before %s",
+			TREE_USED (fndecl) ? "use" : "definition");
+	      parser->oacc_routine = NULL;
+	      return;
+	    }
 
-      /* Process the routine's dimension clauses.  */
-      tree dims = build_oacc_routine_dims (parser->oacc_routine->clauses);
-      replace_oacc_fn_attrib (fndecl, dims);
-      
-      /* Add an "omp declare target" attribute.  */
-      DECL_ATTRIBUTES (fndecl)
-	= tree_cons (get_identifier ("omp declare target"),
-		     NULL_TREE, DECL_ATTRIBUTES (fndecl));
+	  /* Set the routine's level of parallelism.  */
+	  tree dims = build_oacc_routine_dims (parser->oacc_routine->clauses);
+	  replace_oacc_fn_attrib (fndecl, dims);
+
+	  /* Add an "omp declare target" attribute.  */
+	  DECL_ATTRIBUTES (fndecl)
+	    = tree_cons (get_identifier ("omp declare target"),
+			 parser->oacc_routine->clauses,
+			 DECL_ATTRIBUTES (fndecl));
+	}
 
       /* Don't unset parser->oacc_routine here: we may still need it to
 	 diagnose wrong usage.  But, remember that we've used this "#pragma acc
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 1a7c478..8958347 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7073,6 +7073,8 @@  finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	case OMP_CLAUSE_AUTO:
 	case OMP_CLAUSE_INDEPENDENT:
 	case OMP_CLAUSE_SEQ:
+	case OMP_CLAUSE_BIND:
+	case OMP_CLAUSE_NOHOST:
 	  break;
 
 	case OMP_CLAUSE_TILE: