[PR89433] Use 'oacc_verify_routine_clauses' for C/C++ OpenACC 'routine' directives

	gcc/
	PR middle-end/89433
	* omp-general.c (oacc_build_routine_dims): Move some of its
	processing into...
	(oacc_verify_routine_clauses): ... this new function.
	* omp-general.h (oacc_verify_routine_clauses): New prototype.
	gcc/c/
	PR c/89433
	* c-parser.c (c_parser_oacc_routine): Normalize order of clauses.
	(c_finish_oacc_routine): Call oacc_verify_routine_clauses.
	gcc/cp/
	PR c++/89433
	* parser.c (cp_parser_oacc_routine)
	(cp_parser_late_parsing_oacc_routine): Normalize order of clauses.
	(cp_finalize_oacc_routine): Call oacc_verify_routine_clauses.
	gcc/testsuite/
	PR testsuite/89433
	* c-c++-common/goacc/routine-2.c: Update, and move some test
	into...
	* c-c++-common/goacc/routine-level-of-parallelism-1.c: ... this
	new file.

From-SVN: r271344

gcc/ChangeLog

+2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR middle-end/89433
+	* omp-general.c (oacc_build_routine_dims): Move some of its
+	processing into...
+	(oacc_verify_routine_clauses): ... this new function.
+	* omp-general.h (oacc_verify_routine_clauses): New prototype.
+
 2019-05-17  Iain Sandoe  <iain@sandoe.co.uk>
 
 	* config/rs6000/rs6000.c (machopic_output_stub): Adjust the

gcc/c/ChangeLog

 2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
 	PR c/89433
+	* c-parser.c (c_parser_oacc_routine): Normalize order of clauses.
+	(c_finish_oacc_routine): Call oacc_verify_routine_clauses.
+
+	PR c/89433
 	* c-parser.c (c_finish_oacc_routine): Refer to OpenACC 'routine'
 	clauses from "omp declare target" attribute.
 

gcc/c/c-parser.c

@@ -15801,6 +15801,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
       data.clauses
 	= c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
 				     "#pragma acc routine");
+      /* The clauses are in reverse order; fix that to make later diagnostic
+	 emission easier.  */
+      data.clauses = nreverse (data.clauses);
 
       if (TREE_CODE (decl) != FUNCTION_DECL)
 	{
@@ -15815,6 +15818,9 @@ c_parser_oacc_routine (c_parser *parser, enum pragma_context context)
       data.clauses
 	= c_parser_oacc_all_clauses (parser, OACC_ROUTINE_CLAUSE_MASK,
 				     "#pragma acc routine");
+      /* The clauses are in reverse order; fix that to make later diagnostic
+	 emission easier.  */
+      data.clauses = nreverse (data.clauses);
 
       /* Emit a helpful diagnostic if there's another pragma following this
 	 one.  Also don't allow a static assertion declaration, as in the
@@ -15878,6 +15884,8 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl,
       return;
     }
 
+  oacc_verify_routine_clauses (&data->clauses, data->loc);
+
   if (oacc_get_fn_attrib (fndecl))
     {
       error_at (data->loc,

gcc/cp/ChangeLog

 2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
 	PR c++/89433
+	* parser.c (cp_parser_oacc_routine)
+	(cp_parser_late_parsing_oacc_routine): Normalize order of clauses.
+	(cp_finalize_oacc_routine): Call oacc_verify_routine_clauses.
+
+	PR c++/89433
 	* parser.c (cp_finalize_oacc_routine): Refer to OpenACC 'routine'
 	clauses from "omp declare target" attribute.
 

gcc/cp/parser.c

@@ -40136,6 +40136,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
@@ -40232,6 +40235,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.  */
@@ -40266,6 +40272,9 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
 	  return;
 	}
 
+      oacc_verify_routine_clauses (&parser->oacc_routine->clauses,
+				   parser->oacc_routine->loc);
+
       if (oacc_get_fn_attrib (fndecl))
 	{
 	  error_at (parser->oacc_routine->loc,

gcc/omp-general.c

@@ -608,9 +608,61 @@ oacc_set_fn_attrib (tree fn, tree clauses, 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.
+
+   Upon returning, the chain of clauses will contain exactly one clause
+   specifying the level of parallelism.  */
+
+void
+oacc_verify_routine_clauses (tree *clauses, location_t loc)
+{
+  tree c_level = 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),
+		    "... to the previous %qs clause here",
+		    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;
+      default:
+	gcc_unreachable ();
+      }
+  if (c_level == NULL_TREE)
+    {
+      /* Default to an implicit 'seq' clause.  */
+      c_level = build_omp_clause (loc, OMP_CLAUSE_SEQ);
+      OMP_CLAUSE_CHAIN (c_level) = *clauses;
+      *clauses = c_level;
+    }
+}
+
+/*  Process the OpenACC 'routine' directive 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
@@ -632,16 +684,10 @@ oacc_build_routine_dims (tree 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;
 	}
-
-  /* Default to SEQ.  */
-  if (level < 0)
-    level = GOMP_DIM_MAX;
+  gcc_checking_assert (level >= 0);
 
   tree dims = NULL_TREE;
 

gcc/omp-general.h

@@ -84,6 +84,7 @@ extern tree oacc_launch_pack (unsigned code, tree device, unsigned op);
 extern tree oacc_replace_fn_attrib_attr (tree attribs, tree dims);
 extern void oacc_replace_fn_attrib (tree fn, tree dims);
 extern void oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args);
+extern void oacc_verify_routine_clauses (tree *, location_t);
 extern tree oacc_build_routine_dims (tree clauses);
 extern tree oacc_get_fn_attrib (tree fn);
 extern bool offloading_function_p (tree fn);

gcc/testsuite/ChangeLog

 2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
 	PR testsuite/89433
+	* c-c++-common/goacc/routine-2.c: Update, and move some test
+	into...
+	* c-c++-common/goacc/routine-level-of-parallelism-1.c: ... this
+	new file.
+
+	PR testsuite/89433
 	* c-c++-common/goacc/classify-routine.c: Update.
 	* gfortran.dg/goacc/classify-routine.f95: Likewise.
 

gcc/testsuite/c-c++-common/goacc/routine-2.c

-#pragma acc routine gang worker /* { dg-error "multiple loop axes" } */
-void gang (void)
-{
-}
-
-#pragma acc routine worker vector /* { dg-error "multiple loop axes" } */
-void worker (void)
-{
-}
-
-#pragma acc routine vector seq /* { dg-error "multiple loop axes" } */
-void vector (void)
-{
-}
-
-#pragma acc routine seq gang /* { dg-error "multiple loop axes" } */
-void seq (void)
-{
-}
+/* Test invalid use of the OpenACC 'routine' directive.  */
 
 #pragma acc routine (nothing) gang /* { dg-error "not been declared" } */