[OpenACC] More precise diagnostics for 'gang', 'worker', 'vector' clauses with arguments on 'loop' only allowed in 'kernels' regions

Instead of at the location of the 'loop' directive, 'error_at' the location of
the improper clause, and 'inform' at the location of the enclosing parent
compute construct/routine.

The Fortran testcases come with some XFAILing, to be resolved later.

	gcc/
	* omp-low.c (scan_omp_for) <OpenACC>: More precise diagnostics for
	'gang', 'worker', 'vector' clauses with arguments only allowed in
	'kernels' regions.
	gcc/testsuite/
	* c-c++-common/goacc/pr92793-1.c: Extend.
	* gfortran.dg/goacc/pr92793-1.f90: Likewise.

3 files changed, 108 insertions, 10 deletions

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 5392fa7..de5142f 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -2418,30 +2418,39 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
       if (!tgt || is_oacc_parallel_or_serial (tgt))
 	for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
 	  {
-	    char const *check = NULL;
-
+	    tree c_op0;
 	    switch (OMP_CLAUSE_CODE (c))
 	      {
 	      case OMP_CLAUSE_GANG:
-		check = "gang";
+		c_op0 = OMP_CLAUSE_GANG_EXPR (c);
 		break;
 
 	      case OMP_CLAUSE_WORKER:
-		check = "worker";
+		c_op0 = OMP_CLAUSE_WORKER_EXPR (c);
 		break;
 
 	      case OMP_CLAUSE_VECTOR:
-		check = "vector";
+		c_op0 = OMP_CLAUSE_VECTOR_EXPR (c);
 		break;
 
 	      default:
-		break;
+		continue;
 	      }
 
-	    if (check && OMP_CLAUSE_OPERAND (c, 0))
-	      error_at (gimple_location (stmt),
-			"argument not permitted on %qs clause in"
-			" OpenACC %<parallel%> or %<serial%>", check);
+	    if (c_op0)
+	      {
+		error_at (OMP_CLAUSE_LOCATION (c),
+			  "argument not permitted on %qs clause",
+			  omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+		if (tgt)
+		  inform (gimple_location (outer_ctx->stmt),
+			  "enclosing parent compute construct");
+		else if (oacc_get_fn_attrib (current_function_decl))
+		  inform (DECL_SOURCE_LOCATION (current_function_decl),
+			  "enclosing routine");
+		else
+		  gcc_unreachable ();
+	      }
 	  }
 
       if (tgt && is_oacc_kernels (tgt))
diff --git a/gcc/testsuite/c-c++-common/goacc/pr92793-1.c b/gcc/testsuite/c-c++-common/goacc/pr92793-1.c
index d7a2ae4..77ebb20 100644
--- a/gcc/testsuite/c-c++-common/goacc/pr92793-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/pr92793-1.c
@@ -54,3 +54,40 @@ reduction(-:sum  ) /* { dg-line sum2 } */ \
       }
   }
 }
+
+
+void
+a_sl() {
+#pragma acc serial loop /* { dg-message "9: enclosing parent compute construct" } */ \
+    gang(num:5) /* { dg-error "5: argument not permitted on 'gang' clause" } */ \
+  worker(num:5) /* { dg-error "3: argument not permitted on 'worker' clause" } */ \
+   vector(length:5) /* { dg-error "4: argument not permitted on 'vector' clause" } */
+  for (int i = 0; i < 10; i++)
+    ;
+}
+
+void
+a_s_l() {
+#pragma acc serial /* { dg-message "9: enclosing parent compute construct" } */
+  {
+#pragma acc loop \
+       gang(num:5) /* { dg-error "8: argument not permitted on 'gang' clause" } */ \
+   worker(num:5) /* { dg-error "4: argument not permitted on 'worker' clause" } */ \
+  vector(length:5) /* { dg-error "3: argument not permitted on 'vector' clause" } */
+    for (int i = 0; i < 10; i++)
+      ;
+  }
+}
+
+void a_r();
+#pragma acc routine(a_r)
+
+void
+a_r() { /* { dg-message "1: enclosing routine" } */
+#pragma acc loop \
+   gang(num:5) /* { dg-error "4: argument not permitted on 'gang' clause" } */ \
+    worker(num:5) /* { dg-error "5: argument not permitted on 'worker' clause" } */ \
+  vector(length:5) /* { dg-error "3: argument not permitted on 'vector' clause" } */
+  for (int i = 0; i < 10; i++)
+    ;
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr92793-1.f90 b/gcc/testsuite/gfortran.dg/goacc/pr92793-1.f90
index 72dd6b7..23d6886 100644
--- a/gcc/testsuite/gfortran.dg/goacc/pr92793-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/pr92793-1.f90
@@ -45,3 +45,55 @@ subroutine check ()
   end do
 !$acc end  parallel
 end subroutine check
+
+
+subroutine gwv_sl ()
+  implicit none (type, external)
+  integer :: i
+
+  !$acc serial loop &
+  !$acc &       gang(num:5) & ! { dg-error "argument not permitted on 'gang' clause" "TODO" { xfail *-*-* } }
+  !$acc & & ! { dg-bogus "14: argument not permitted on 'gang' clause" "TODO" { xfail *-*-* } .+6 }
+  !$acc &    worker(num:5) & ! { dg-error "argument not permitted on 'worker' clause" "TODO" { xfail *-*-* } }
+  !$acc & & ! { dg-bogus "14: argument not permitted on 'worker' clause" "TODO" { xfail *-*-* } .+4 }
+  !$acc &     vector(length:5) & ! { dg-error "argument not permitted on 'vector' clause" "TODO" { xfail *-*-* } }
+  !$acc & ! { dg-bogus "14: argument not permitted on 'vector' clause" "TODO" { xfail *-*-* } .+2 }
+  ! { dg-message "99: enclosing parent compute construct" "" { target *-*-* } .-1 }
+  do i = 0, 10
+  end do
+  !$acc end serial loop
+end subroutine gwv_sl
+
+subroutine gwv_s_l ()
+  implicit none (type, external)
+  integer :: i
+
+  !$acc serial ! { dg-message "72: enclosing parent compute construct" }
+  !$acc loop &
+  !$acc &         gang(num:5) & ! { dg-error "argument not permitted on 'gang' clause" "TODO" { xfail *-*-* } }
+  !$acc & & ! { dg-bogus "14: argument not permitted on 'gang' clause" "TODO" { xfail *-*-* } .+5 }
+  !$acc &   worker(num:5) & ! { dg-error "argument not permitted on 'worker' clause" "TODO" { xfail *-*-* } }
+  !$acc & & ! { dg-bogus "14: argument not permitted on 'worker' clause" "TODO" { xfail *-*-* } .+3 }
+  !$acc &      vector(length:5) & ! { dg-error "argument not permitted on 'vector' clause" "TODO" { xfail *-*-* } }
+  !$acc & ! { dg-bogus "14: argument not permitted on 'vector' clause" "TODO" { xfail *-*-* } .+1 }
+  do i = 0, 10
+  end do
+  !$acc end serial
+end subroutine gwv_s_l
+
+subroutine gwv_r () ! { dg-message "16: enclosing routine" }
+  implicit none (type, external)
+  integer :: i
+
+  !$acc routine(gwv_r)
+
+  !$acc loop &
+  !$acc &     gang(num:5) & ! { dg-error "argument not permitted on 'gang' clause" "TODO" { xfail *-*-* } }
+  !$acc & & ! { dg-bogus "14: argument not permitted on 'gang' clause" "TODO" { xfail *-*-* } .+5 }
+  !$acc &      worker(num:5) & ! { dg-error "argument not permitted on 'worker' clause" "TODO" { xfail *-*-* } }
+  !$acc & & ! { dg-bogus "14: argument not permitted on 'worker' clause" "TODO" { xfail *-*-* } .+3 }
+  !$acc &    vector(length:5) & ! { dg-error "argument not permitted on 'vector' clause" "TODO" { xfail *-*-* } }
+  !$acc & ! { dg-bogus "14: argument not permitted on 'vector' clause" "TODO" { xfail *-*-* } .+1 }
+  do i = 0, 10
+  end do
+end subroutine gwv_r