[PR89433] Repeated use of the C/C++ OpenACC 'routine' directive

	gcc/
	PR middle-end/89433
	* omp-general.c (oacc_verify_routine_clauses): Change formal
	parameters.  Add checking if already marked with an OpenACC
	'routine' directive.  Adjust all users.
	gcc/c/
	PR c/89433
	* c-parser.c (c_finish_oacc_routine): Rework checking if already
	marked with an OpenACC 'routine' directive.
	gcc/cp/
	PR c++/89433
	* parser.c (cp_finalize_oacc_routine): Rework checking if already
	marked with an OpenACC 'routine' directive.
	gcc/testsuite/
	PR testsuite/89433
	* c-c++-common/goacc/routine-5.c: Update.
	* c-c++-common/goacc/routine-level-of-parallelism-1.c: Likewise.
	* c-c++-common/goacc/routine-level-of-parallelism-2.c: New file.

From-SVN: r271345

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index b8a4358..bee1292 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,6 +1,11 @@
 2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
        PR middle-end/89433
+       * omp-general.c (oacc_verify_routine_clauses): Change formal
+       parameters.  Add checking if already marked with an OpenACC
+       'routine' directive.  Adjust all users.
+
+       PR middle-end/89433
        * omp-general.c (oacc_build_routine_dims): Move some of its
        processing into...
        (oacc_verify_routine_clauses): ... this new function.

diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 1393e8f..cfbd164 100644 (file)
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,6 +1,10 @@
 2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
        PR c/89433
+       * c-parser.c (c_finish_oacc_routine): Rework checking if already
+       marked with an OpenACC 'routine' directive.
+
+       PR c/89433
        * c-parser.c (c_parser_oacc_routine): Normalize order of clauses.
        (c_finish_oacc_routine): Call oacc_verify_routine_clauses.
 

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 8337f1c..8f61024 100644 (file)
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15884,35 +15884,39 @@ 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))
+  int compatible
+    = oacc_verify_routine_clauses (fndecl, &data->clauses, data->loc,
+                                  "#pragma acc routine");
+  if (compatible < 0)
     {
-      error_at (data->loc,
-               "%<#pragma acc routine%> already applied to %qD", fndecl);
       data->error_seen = true;
       return;
     }
-
-  if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+  if (compatible > 0)
     {
-      error_at (data->loc,
-               TREE_USED (fndecl)
-               ? G_("%<#pragma acc routine%> must be applied before use")
-               : G_("%<#pragma acc routine%> must be applied before "
-                    "definition"));
-      data->error_seen = true;
-      return;
     }
+  else
+    {
+      if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+       {
+         error_at (data->loc,
+                   TREE_USED (fndecl)
+                   ? G_("%<#pragma acc routine%> must be applied before use")
+                   : G_("%<#pragma acc routine%> must be applied before"
+                        " definition"));
+         data->error_seen = true;
+         return;
+       }
 
-  /* Process the routine's dimension clauses.  */
-  tree dims = oacc_build_routine_dims (data->clauses);
-  oacc_replace_fn_attrib (fndecl, dims);
+      /* Set the routine's level of parallelism.  */
+      tree dims = oacc_build_routine_dims (data->clauses);
+      oacc_replace_fn_attrib (fndecl, dims);
 
-  /* Add an "omp declare target" attribute.  */
-  DECL_ATTRIBUTES (fndecl)
-    = tree_cons (get_identifier ("omp declare target"),
-                data->clauses, DECL_ATTRIBUTES (fndecl));
+      /* Add an "omp declare target" attribute.  */
+      DECL_ATTRIBUTES (fndecl)
+       = tree_cons (get_identifier ("omp declare target"),
+                    data->clauses, DECL_ATTRIBUTES (fndecl));
+    }
 
   /* Remember that we've used this "#pragma acc routine".  */
   data->fndecl_seen = true;

diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index 2f1e06c..39aaab9 100644 (file)
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,6 +1,10 @@
 2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
        PR c++/89433
+       * parser.c (cp_finalize_oacc_routine): Rework checking if already
+       marked with an OpenACC 'routine' directive.
+
+       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.

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index aa6507e..6705d64 100644 (file)
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40272,36 +40272,42 @@ 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))
+      int compatible
+       = oacc_verify_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%> already applied to %qD", fndecl);
          parser->oacc_routine = NULL;
          return;
        }
-
-      if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+      if (compatible > 0)
        {
-         error_at (parser->oacc_routine->loc,
-                   TREE_USED (fndecl)
-                   ? G_("%<#pragma acc routine%> must be applied before use")
-                   : G_("%<#pragma acc routine%> must be applied before "
-                        "definition"));
-         parser->oacc_routine = NULL;
-         return;
        }
+      else
+       {
+         if (TREE_USED (fndecl) || (!is_defn && DECL_SAVED_TREE (fndecl)))
+           {
+             error_at (parser->oacc_routine->loc,
+                       TREE_USED (fndecl)
+                       ? G_("%<#pragma acc routine%> must be applied before"
+                            " use")
+                       : G_("%<#pragma acc routine%> must be applied before"
+                            " definition"));
+             parser->oacc_routine = NULL;
+             return;
+           }
 
-      /* Process the routine's dimension clauses.  */
-      tree dims = oacc_build_routine_dims (parser->oacc_routine->clauses);
-      oacc_replace_fn_attrib (fndecl, dims);
+         /* Set the routine's level of parallelism.  */
+         tree dims = oacc_build_routine_dims (parser->oacc_routine->clauses);
+         oacc_replace_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));
+         /* 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/omp-general.c b/gcc/omp-general.c
index f1d859b..82f0a04 100644 (file)
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -610,11 +610,14 @@ oacc_set_fn_attrib (tree fn, tree clauses, vec<tree> *args)
 
 /* Verify OpenACC routine clauses.
 
+   Returns 0 if FNDECL should be marked with an OpenACC 'routine' directive, 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.  */
 
-void
-oacc_verify_routine_clauses (tree *clauses, location_t loc)
+int
+oacc_verify_routine_clauses (tree fndecl, tree *clauses, location_t loc,
+                            const char *routine_str)
 {
   tree c_level = NULL_TREE;
   tree c_p = NULL_TREE;
@@ -659,6 +662,80 @@ oacc_verify_routine_clauses (tree *clauses, location_t loc)
       OMP_CLAUSE_CHAIN (c_level) = *clauses;
       *clauses = c_level;
     }
+  /* In *clauses, we now have exactly one clause specifying the level of
+     parallelism.  */
+
+  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;
+      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;
+         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;
+       }
+      /* 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 with an OpenACC 'routine' directive",
+                 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 with an OpenACC 'routine' directive",
+                 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),
+               "... with %qs clause here",
+               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, "... without %qs clause near to here",
+                 omp_clause_code_name[OMP_CLAUSE_CODE (c_diag)]);
+       }
+      /* Incompatible.  */
+      return -1;
+    }
+
+  return 0;
 }
 
 /*  Process the OpenACC 'routine' directive clauses to generate an attribute

diff --git a/gcc/omp-general.h b/gcc/omp-general.h
index 4241c33..f96d3c7 100644 (file)
--- a/gcc/omp-general.h
+++ b/gcc/omp-general.h
@@ -84,7 +84,8 @@ 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 int oacc_verify_routine_clauses (tree, tree *, location_t,
+                                       const char *);
 extern tree oacc_build_routine_dims (tree clauses);
 extern tree oacc_get_fn_attrib (tree fn);
 extern bool offloading_function_p (tree fn);

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 4b07888..ed94f19 100644 (file)
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,6 +1,11 @@
 2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
 
        PR testsuite/89433
+       * c-c++-common/goacc/routine-5.c: Update.
+       * c-c++-common/goacc/routine-level-of-parallelism-1.c: Likewise.
+       * c-c++-common/goacc/routine-level-of-parallelism-2.c: New file.
+
+       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

diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c
index b967a74..a68c6be 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -150,61 +150,19 @@ void f_static_assert();
 
 #pragma acc routine
 __extension__ extern void ex1();
-#pragma acc routine (ex1) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex1" } */
+#pragma acc routine (ex1) worker /* { dg-error "has already been marked with an OpenACC 'routine' directive" } */
 
 #pragma acc routine
 __extension__ __extension__ __extension__ __extension__ __extension__ void ex2()
 {
 }
-#pragma acc routine (ex2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex2" } */
+#pragma acc routine (ex2) worker /* { dg-error "has already been marked with an OpenACC 'routine' directive" } */
 
 #pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
 __extension__ int ex3;
 #pragma acc routine (ex3) /* { dg-error ".ex3. does not refer to a function" } */
 
 
-/* "#pragma acc routine" already applied.  */
-
-extern void fungsi_1();
-#pragma acc routine(fungsi_1) gang
-#pragma acc routine(fungsi_1) gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
-#pragma acc routine(fungsi_1) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
-#pragma acc routine(fungsi_1) vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
-
-#pragma acc routine seq
-extern void fungsi_2();
-#pragma acc routine(fungsi_2) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
-#pragma acc routine(fungsi_2) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
-#pragma acc routine(fungsi_2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
-
-#pragma acc routine vector
-extern void fungsi_3();
-#pragma acc routine vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_3." } */
-void fungsi_3()
-{
-}
-
-extern void fungsi_4();
-#pragma acc routine (fungsi_4) worker
-#pragma acc routine gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_4." } */
-void fungsi_4()
-{
-}
-
-#pragma acc routine gang
-void fungsi_5()
-{
-}
-#pragma acc routine (fungsi_5) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_5." } */
-
-#pragma acc routine seq
-void fungsi_6()
-{
-}
-#pragma acc routine seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_6." } */
-extern void fungsi_6();
-
-
 /* "#pragma acc routine" must be applied before.  */
 
 void Bar ();

diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
index ab0414b..4fdeb14 100644 (file)
--- a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-1.c
@@ -41,10 +41,10 @@ void s_2 (void)
 void g_3 (void)
 {
 }
-#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_3." } */ \
+#pragma acc routine (g_3) \
   gang \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
-#pragma acc routine (g_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_3." } */ \
+#pragma acc routine (g_3) \
   gang \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
 
@@ -52,10 +52,10 @@ extern void w_3 (void);
 #pragma acc routine (w_3) \
   worker \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
-#pragma acc routine (w_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_3." } */ \
+#pragma acc routine (w_3) \
   worker \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
-#pragma acc routine (w_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_3." } */ \
+#pragma acc routine (w_3) \
   worker \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
 
@@ -65,10 +65,10 @@ extern void w_3 (void);
 void v_3 (void)
 {
 }
-#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_3." } */ \
+#pragma acc routine (v_3) \
   vector \
   worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
-#pragma acc routine (v_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_3." } */ \
+#pragma acc routine (v_3) \
   vector \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
 
@@ -76,10 +76,10 @@ extern void s_3 (void);
 #pragma acc routine (s_3) \
   seq \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
-#pragma acc routine (s_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_3." } */ \
+#pragma acc routine (s_3) \
   seq \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */
-#pragma acc routine (s_3) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_3." } */ \
+#pragma acc routine (s_3) \
   seq \
   worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
 
@@ -90,12 +90,12 @@ extern void s_3 (void);
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
 extern void g_4 (void);
-#pragma acc routine (g_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_4." } */ \
+#pragma acc routine (g_4) \
   gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
   worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */
-#pragma acc routine (g_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_4." } */ \
+#pragma acc routine (g_4) \
   gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
@@ -107,12 +107,12 @@ extern void w_4 (void);
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
-#pragma acc routine (w_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_4." } */ \
+#pragma acc routine (w_4) \
   worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
-#pragma acc routine (w_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_4." } */ \
+#pragma acc routine (w_4) \
   worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
@@ -126,12 +126,12 @@ extern void w_4 (void);
 void v_4 (void)
 {
 }
-#pragma acc routine (v_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_4." } */ \
+#pragma acc routine (v_4) \
   vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
   worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */
-#pragma acc routine (v_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_4." } */ \
+#pragma acc routine (v_4) \
   vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
   seq /* { dg-error ".seq. specifies a conflicting level of parallelism" } */ \
@@ -145,12 +145,12 @@ void v_4 (void)
 void s_4 (void)
 {
 }
-#pragma acc routine (s_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_4." } */ \
+#pragma acc routine (s_4) \
   seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
   gang /* { dg-error ".gang. specifies a conflicting level of parallelism" } */ \
   worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */
-#pragma acc routine (s_4) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_4." } */ \
+#pragma acc routine (s_4) \
   seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
   worker /* { dg-error ".worker. specifies a conflicting level of parallelism" } */ \
   vector /* { dg-error ".vector. specifies a conflicting level of parallelism" } */ \
@@ -168,7 +168,7 @@ void s_4 (void)
 void g_5 (void)
 {
 }
-#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \
+#pragma acc routine (g_5) \
   gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
   vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
   /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
@@ -176,7 +176,7 @@ void g_5 (void)
   /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
   /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
-#pragma acc routine (g_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*g_5." } */ \
+#pragma acc routine (g_5) \
   gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
   seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
   /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
@@ -194,7 +194,7 @@ void g_5 (void)
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
 extern void w_5 (void);
-#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_5." } */ \
+#pragma acc routine (w_5) \
   worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
   vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
   /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
@@ -202,7 +202,7 @@ extern void w_5 (void);
   /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
-#pragma acc routine (w_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*w_5." } */ \
+#pragma acc routine (w_5) \
   worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
   seq seq /* { dg-error "too many 'seq' clauses" } */ \
   /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
@@ -220,7 +220,7 @@ extern void w_5 (void);
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
 extern void v_5 (void);
-#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_5." } */ \
+#pragma acc routine (v_5) \
   vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
   seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
   /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
@@ -228,7 +228,7 @@ extern void v_5 (void);
   /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
-#pragma acc routine (v_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*v_5." } */ \
+#pragma acc routine (v_5) \
   vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
@@ -246,7 +246,7 @@ extern void s_5 (void);
   /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
-#pragma acc routine (s_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_5." } */ \
+#pragma acc routine (s_5) \
   seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
   vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
   /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
@@ -254,7 +254,7 @@ extern void s_5 (void);
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
   /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
-#pragma acc routine (s_5) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*s_5." } */ \
+#pragma acc routine (s_5) \
   seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
   worker worker /* { dg-error "too many 'worker' clauses" } */ \
   /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
@@ -262,3 +262,188 @@ extern void s_5 (void);
   /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
   gang gang /* { dg-error "too many 'gang' clauses" } */ \
   /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+
+/* Like the *_5 tests, but with the order of clauses changed in the second and
+   following routine directives for the specific *_5 function.  */
+
+#pragma acc routine \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+void g_6 (void)
+{
+}
+#pragma acc routine (g_6) \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*g_6\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" "" { target *-*-* } .-1 } */ \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#pragma acc routine (g_6) \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_6\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" "" { target *-*-* } .-1 } */ \
+  gang gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+#pragma acc routine \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  vector vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+extern void w_6 (void);
+#pragma acc routine (w_6) \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*w_6\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" "" { target *-*-* } .-1 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#pragma acc routine (w_6) \
+  seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_6\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" "" { target *-*-* } .-1 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+#pragma acc routine \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+extern void v_6 (void);
+#pragma acc routine (v_6) \
+  seq seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_6\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" "" { target *-*-* } .-1 } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#pragma acc routine (v_6) \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*v_6\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" "" { target *-*-* } .-1 } */ \
+  vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  seq seq seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+extern void s_6 (void);
+#pragma acc routine (s_6) \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#pragma acc routine (s_6) \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_6\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" "" { target *-*-* } .-1 } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  worker worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error ".worker. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+#pragma acc routine (s_6) \
+  worker worker /* { dg-error "too many 'worker' clauses" } */ \
+  /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_6\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" "" { target *-*-* } .-1 } */ \
+  seq seq seq /* { dg-error "too many 'seq' clauses" } */ \
+  /* { dg-error ".seq. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  vector vector vector vector /* { dg-error "too many 'vector' clauses" } */ \
+  /* { dg-error ".vector. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */ \
+  gang gang /* { dg-error "too many 'gang' clauses" } */ \
+  /* { dg-error ".gang. specifies a conflicting level of parallelism" "" { target *-*-* } .-1 } */
+
+
+/* Like the *_6 tests, but without all the duplicate clauses, so that the
+   routine directives are valid in isolation.  */
+
+#pragma acc routine \
+  gang
+void g_7 (void)
+{
+}
+#pragma acc routine (g_7) \
+  vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*g_7\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+#pragma acc routine (g_7) \
+  seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_7\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+
+#pragma acc routine \
+  worker
+extern void w_7 (void);
+#pragma acc routine (w_7) \
+  vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*w_7\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+#pragma acc routine (w_7) \
+  seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_7\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+
+#pragma acc routine \
+  vector
+extern void v_7 (void);
+#pragma acc routine (v_7) \
+  seq /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_7\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+#pragma acc routine (v_7) \
+  gang /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*v_7\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+
+extern void s_7 (void);
+#pragma acc routine (s_7) \
+  seq
+#pragma acc routine (s_7) \
+  vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_7\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+#pragma acc routine (s_7) \
+  worker /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_7\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+
+
+/* Test cases for implicit seq clause.  */
+
+#pragma acc routine \
+  gang
+void g_8 (void)
+{
+}
+#pragma acc routine (g_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*g_8\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+
+#pragma acc routine \
+  worker
+extern void w_8 (void);
+#pragma acc routine (w_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*w_8\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+
+#pragma acc routine \
+  vector
+extern void v_8 (void);
+#pragma acc routine (v_8) /* { dg-error "incompatible .seq. clause when applying .#pragma acc routine. to .\[void \]*v_8\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+
+extern void s_8 (void);
+#pragma acc routine (s_8)
+#pragma acc routine (s_8) \
+  vector /* { dg-error "incompatible .vector. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+#pragma acc routine (s_8) \
+  gang /* { dg-error "incompatible .gang. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */
+#pragma acc routine (s_8) \
+  worker /* { dg-error "incompatible .worker. clause when applying .#pragma acc routine. to .\[void \]*s_8\[\\(\\)\]*., which has already been marked with an OpenACC 'routine' directive" } */

diff --git a/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
new file mode 100644 (file)
index 0000000..a066f2b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/routine-level-of-parallelism-2.c
@@ -0,0 +1,71 @@
+/* Test various aspects of clauses specifying compatible levels of parallelism
+   with the OpenACC 'routine' directive.  The Fortran counterpart is
+   '../../gfortran.dg/goacc/routine-level-of-parallelism-1.f90'.  */
+
+#pragma acc routine gang
+void g_1 (void)
+{
+}
+#pragma acc routine (g_1) gang
+#pragma acc routine (g_1) gang
+
+
+extern void w_1 (void);
+#pragma acc routine (w_1) worker
+#pragma acc routine (w_1) worker
+#pragma acc routine (w_1) worker
+
+
+#pragma acc routine vector
+extern void v_1 (void);
+#pragma acc routine (v_1) vector
+#pragma acc routine (v_1) vector
+
+
+/* Also test the implicit seq clause.  */
+
+#pragma acc routine seq
+extern void s_1_1 (void);
+#pragma acc routine (s_1_1)
+#pragma acc routine (s_1_1) seq
+#pragma acc routine (s_1_1)
+#pragma acc routine (s_1_1) seq
+
+#pragma acc routine
+extern void s_1_2 (void);
+#pragma acc routine (s_1_2)
+#pragma acc routine (s_1_2) seq
+#pragma acc routine (s_1_2)
+#pragma acc routine (s_1_2) seq
+
+extern void s_2_1 (void);
+#pragma acc routine (s_2_1) seq
+#pragma acc routine (s_2_1)
+#pragma acc routine (s_2_1) seq
+#pragma acc routine (s_2_1)
+#pragma acc routine (s_2_1) seq
+
+extern void s_2_2 (void);
+#pragma acc routine (s_2_2)
+#pragma acc routine (s_2_2)
+#pragma acc routine (s_2_2) seq
+#pragma acc routine (s_2_2)
+#pragma acc routine (s_2_2) seq
+
+#pragma acc routine seq
+void s_3_1 (void)
+{
+}
+#pragma acc routine (s_3_1)
+#pragma acc routine (s_3_1) seq
+#pragma acc routine (s_3_1)
+#pragma acc routine (s_3_1) seq
+
+#pragma acc routine
+void s_3_2 (void)
+{
+}
+#pragma acc routine (s_3_2)
+#pragma acc routine (s_3_2) seq
+#pragma acc routine (s_3_2)
+#pragma acc routine (s_3_2) seq

diff --git a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90 b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
index 75dd1b0..83b8c24 100644 (file)
--- a/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/routine-level-of-parallelism-1.f90
@@ -1,6 +1,6 @@
-! Test various aspects of clauses specifying compatible levels of
-! parallelism with the OpenACC routine directive.  The Fortran counterpart is
-! c-c++-common/goacc/routine-level-of-parallelism-2.c
+! Test various aspects of clauses specifying compatible levels of parallelism
+! with the OpenACC routine directive.  The C/C++ counterpart is
+! '../../c-c++-common/goacc/routine-level-of-parallelism-2.c'.
 
 subroutine g_1
   !$acc routine gang