public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses
@ 2017-05-11 12:28 Thomas Schwinge
  2017-05-14 10:40 ` Thomas Schwinge
                   ` (2 more replies)
  0 siblings, 3 replies; 5+ messages in thread
From: Thomas Schwinge @ 2017-05-11 12:28 UTC (permalink / raw)
  To: gcc-patches, Jakub Jelinek

Hi!

Building on the other pending patches (I'll soon commit the approved
ones), we can then support the num_gangs, num_workers, vector_length
clauses for the OpenACC 2.5 kernels construct.  OK for trunk?

commit a689c52cde71960bc08ae30c3f88980f66fdd0b8
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Thu May 11 13:43:28 2017 +0200

    OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses
    
            gcc/c/
            * c-parser.c (OACC_KERNELS_CLAUSE_MASK): Add
            "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
            "VECTOR_LENGTH".
            gcc/cp/
            * parser.c (OACC_KERNELS_CLAUSE_MASK): Add
            "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
            "VECTOR_LENGTH".
            gcc/fortran/
            * openmp.c (OACC_KERNELS_CLAUSES): Add "OMP_CLAUSE_NUM_GANGS",
            "OMP_CLAUSE_NUM_WORKERS", "OMP_CLAUSE_VECTOR_LENGTH".
            gcc/
            * omp-offload.c (execute_oacc_device_lower): Remove the
            parallelism dimensions function attributes for unparallelized
            OpenACC kernels constructs.
            gcc/testsuite/
            * c-c++-common/goacc/parallel-dims-1.c: Update.
            * c-c++-common/goacc/parallel-dims-2.c: Likewise.
            * c-c++-common/goacc/routine-1.c: Likewise.
            * c-c++-common/goacc/uninit-dim-clause.c: Likewise.
            * g++.dg/goacc/template.C: Likewise.
            * gfortran.dg/goacc/kernels-tree.f95: Likewise.
            * gfortran.dg/goacc/routine-3.f90: Likewise.
            * gfortran.dg/goacc/sie.f95: Likewise.
            * gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
            libgomp/
            * testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: Update.
            * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
            * testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise.
---
 gcc/c/c-parser.c                                   |   3 +
 gcc/cp/parser.c                                    |   3 +
 gcc/fortran/openmp.c                               |   3 +-
 gcc/omp-offload.c                                  |   9 ++
 gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c |   3 +
 gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c | 152 +++++++++++++++++++--
 gcc/testsuite/c-c++-common/goacc/routine-1.c       |   7 +
 .../c-c++-common/goacc/uninit-dim-clause.c         |  20 ++-
 gcc/testsuite/g++.dg/goacc/template.C              |   4 +
 gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95   |   6 +-
 gcc/testsuite/gfortran.dg/goacc/routine-3.f90      |   6 +
 gcc/testsuite/gfortran.dg/goacc/sie.f95            |  86 +++++++++++-
 .../gfortran.dg/goacc/uninit-dim-clause.f95        |  18 ++-
 .../libgomp.oacc-c-c++-common/kernels-loop-2.c     |  21 ++-
 .../libgomp.oacc-c-c++-common/parallel-dims.c      |  35 +++++
 .../libgomp.oacc-fortran/kernels-loop-2.f95        |  13 +-
 16 files changed, 358 insertions(+), 31 deletions(-)

diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 90d2d17..c0d733c 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -13978,11 +13978,14 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
 #define OACC_PARALLEL_CLAUSE_MASK					\
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 17d2679..0578e81 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -36438,11 +36438,14 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
 #define OACC_PARALLEL_CLAUSE_MASK					\
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index 89eecfa..7b18a1d 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -1926,7 +1926,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
    | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE \
    | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
 #define OACC_KERNELS_CLAUSES \
-  (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_DEVICEPTR	      \
+  (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS	      \
+   | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
    | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT		      \
    | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY      \
    | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT	      \
diff --git gcc/omp-offload.c gcc/omp-offload.c
index e954ee9..fab26d0 100644
--- gcc/omp-offload.c
+++ gcc/omp-offload.c
@@ -1452,6 +1452,15 @@ execute_oacc_device_lower ()
     = (lookup_attribute ("oacc kernels parallelized",
 			 DECL_ATTRIBUTES (current_function_decl)) != NULL);
 
+  /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1
+     kernels, so remove the parallelism dimensions function attributes
+     potentially set earlier on.  */
+  if (is_oacc_kernels && !is_oacc_kernels_parallelized)
+    {
+      oacc_set_fn_attrib (current_function_decl, NULL, NULL);
+      attrs = oacc_get_fn_attrib (current_function_decl);
+    }
+
   /* Discover, partition and process the loops.  */
   oacc_loop *loops = oacc_loop_discovery ();
   int fn_level = oacc_fn_attrib_level (attrs);
diff --git gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
index a85d3d3..57f682f 100644
--- gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
+++ gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
@@ -3,6 +3,9 @@
 
 void f(int i)
 {
+#pragma acc kernels num_gangs(i) num_workers(i) vector_length(i)
+  ;
+
 #pragma acc parallel num_gangs(i) num_workers(i) vector_length(i)
   ;
 }
diff --git gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
index 30a3d17..acfbe7f 100644
--- gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
+++ gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
@@ -1,18 +1,15 @@
 /* Invalid use of OpenACC parallelism dimensions clauses: num_gangs,
    num_workers, vector_length.  */
 
-void acc_kernels(int i)
+void f(int i, float f)
 {
-#pragma acc kernels num_gangs(i) /* { dg-error "'num_gangs' is not valid for '#pragma acc kernels'" } */
+#pragma acc kernels num_gangs /* { dg-error "expected '\\(' before end of line" } */
   ;
-#pragma acc kernels num_workers(i) /* { dg-error "'num_workers' is not valid for '#pragma acc kernels'" } */
+#pragma acc kernels num_workers /* { dg-error "expected '\\(' before end of line" } */
   ;
-#pragma acc kernels vector_length(i) /* { dg-error "'vector_length' is not valid for '#pragma acc kernels'" } */
+#pragma acc kernels vector_length /* { dg-error "expected '\\(' before end of line" } */
   ;
-}
 
-void acc_parallel(int i, float f)
-{
 #pragma acc parallel num_gangs /* { dg-error "expected '\\(' before end of line" } */
   ;
 #pragma acc parallel num_workers /* { dg-error "expected '\\(' before end of line" } */
@@ -20,6 +17,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length /* { dg-error "expected '\\(' before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs( /* { dg-error "expected (primary-|)expression before end of line" } */
+  ;
+#pragma acc kernels num_workers( /* { dg-error "expected (primary-|)expression before end of line" } */
+  ;
+#pragma acc kernels vector_length( /* { dg-error "expected (primary-|)expression before end of line" } */
+  ;
+
 #pragma acc parallel num_gangs( /* { dg-error "expected (primary-|)expression before end of line" } */
   ;
 #pragma acc parallel num_workers( /* { dg-error "expected (primary-|)expression before end of line" } */
@@ -27,6 +32,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length( /* { dg-error "expected (primary-|)expression before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
+  ;
+#pragma acc kernels num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
+  ;
+#pragma acc kernels vector_length() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
+  ;
+
 #pragma acc parallel num_gangs() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
   ;
 #pragma acc parallel num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
@@ -34,6 +47,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
   ;
 
+
+#pragma acc kernels num_gangs(1 /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels num_workers(1 /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels vector_length(1 /* { dg-error "expected '\\)' before end of line" } */
+  ;
+
 #pragma acc parallel num_gangs(1 /* { dg-error "expected '\\)' before end of line" } */
   ;
 #pragma acc parallel num_workers(1 /* { dg-error "expected '\\)' before end of line" } */
@@ -41,6 +62,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1 /* { dg-error "expected '\\)' before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs(i /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels num_workers(i /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels vector_length(i /* { dg-error "expected '\\)' before end of line" } */
+  ;
+
 #pragma acc parallel num_gangs(i /* { dg-error "expected '\\)' before end of line" } */
   ;
 #pragma acc parallel num_workers(i /* { dg-error "expected '\\)' before end of line" } */
@@ -48,6 +77,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(i /* { dg-error "expected '\\)' before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs(1 i /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels num_workers(1 i /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels vector_length(1 i /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+
 #pragma acc parallel num_gangs(1 i /* { dg-error "expected '\\)' before 'i'" } */
   ;
 #pragma acc parallel num_workers(1 i /* { dg-error "expected '\\)' before 'i'" } */
@@ -55,6 +92,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1 i /* { dg-error "expected '\\)' before 'i'" } */
   ;
 
+
+#pragma acc kernels num_gangs(1 i) /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels num_workers(1 i) /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels vector_length(1 i) /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+
 #pragma acc parallel num_gangs(1 i) /* { dg-error "expected '\\)' before 'i'" } */
   ;
 #pragma acc parallel num_workers(1 i) /* { dg-error "expected '\\)' before 'i'" } */
@@ -62,6 +107,17 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1 i) /* { dg-error "expected '\\)' before 'i'" } */
   ;
 
+
+#pragma acc kernels num_gangs(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
+  ;
+#pragma acc kernels num_workers(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
+  ;
+#pragma acc kernels vector_length(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
+  ;
+
 #pragma acc parallel num_gangs(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
   /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
   ;
@@ -72,6 +128,14 @@ void acc_parallel(int i, float f)
   /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
   ;
 
+
+#pragma acc kernels num_gangs(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  ;
+#pragma acc kernels num_workers(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  ;
+#pragma acc kernels vector_length(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  ;
+
 #pragma acc parallel num_gangs(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
   ;
 #pragma acc parallel num_workers(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
@@ -79,11 +143,27 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
   ;
 
-#pragma acc parallel num_gangs(num_gangs) /* { dg-error "'num_gangs' (un|was not )declared" } */
+
+#pragma acc kernels num_gangs(num_gangs_k) /* { dg-error "'num_gangs_k' (un|was not )declared" } */
   ;
-#pragma acc parallel num_workers(num_workers) /* { dg-error "'num_workers' (un|was not )declared" } */
+#pragma acc kernels num_workers(num_workers_k) /* { dg-error "'num_workers_k' (un|was not )declared" } */
   ;
-#pragma acc parallel vector_length(vector_length) /* { dg-error "'vector_length' (un|was not )declared" } */
+#pragma acc kernels vector_length(vector_length_k) /* { dg-error "'vector_length_k' (un|was not )declared" } */
+  ;
+
+#pragma acc parallel num_gangs(num_gangs_p) /* { dg-error "'num_gangs_p' (un|was not )declared" } */
+  ;
+#pragma acc parallel num_workers(num_workers_p) /* { dg-error "'num_workers_p' (un|was not )declared" } */
+  ;
+#pragma acc parallel vector_length(vector_length_p) /* { dg-error "'vector_length_p' (un|was not )declared" } */
+  ;
+
+
+#pragma acc kernels num_gangs(f) /* { dg-error "'num_gangs' expression must be integral" } */
+  ;
+#pragma acc kernels num_workers(f) /* { dg-error "'num_workers' expression must be integral" } */
+  ;
+#pragma acc kernels vector_length(f) /* { dg-error "'vector_length' expression must be integral" } */
   ;
 
 #pragma acc parallel num_gangs(f) /* { dg-error "'num_gangs' expression must be integral" } */
@@ -93,6 +173,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(f) /* { dg-error "'vector_length' expression must be integral" } */
   ;
 
+
+#pragma acc kernels num_gangs((float) 1) /* { dg-error "'num_gangs' expression must be integral" } */
+  ;
+#pragma acc kernels num_workers((float) 1) /* { dg-error "'num_workers' expression must be integral" } */
+  ;
+#pragma acc kernels vector_length((float) 1) /* { dg-error "'vector_length' expression must be integral" } */
+  ;
+
 #pragma acc parallel num_gangs((float) 1) /* { dg-error "'num_gangs' expression must be integral" } */
   ;
 #pragma acc parallel num_workers((float) 1) /* { dg-error "'num_workers' expression must be integral" } */
@@ -100,6 +188,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length((float) 1) /* { dg-error "'vector_length' expression must be integral" } */
   ;
 
+
+#pragma acc kernels num_gangs(0) /* { dg-warning "'num_gangs' value must be positive" } */
+  ;
+#pragma acc kernels num_workers(0) /* { dg-warning "'num_workers' value must be positive" } */
+  ;
+#pragma acc kernels vector_length(0) /* { dg-warning "'vector_length' value must be positive" } */
+  ;
+
 #pragma acc parallel num_gangs(0) /* { dg-warning "'num_gangs' value must be positive" } */
   ;
 #pragma acc parallel num_workers(0) /* { dg-warning "'num_workers' value must be positive" } */
@@ -107,6 +203,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(0) /* { dg-warning "'vector_length' value must be positive" } */
   ;
 
+
+#pragma acc kernels num_gangs((int) -1.2) /* { dg-warning "'num_gangs' value must be positive" } */
+  ;
+#pragma acc kernels num_workers((int) -1.2) /* { dg-warning "'num_workers' value must be positive" } */
+  ;
+#pragma acc kernels vector_length((int) -1.2) /* { dg-warning "'vector_length' value must be positive" } */
+  ;
+
 #pragma acc parallel num_gangs((int) -1.2) /* { dg-warning "'num_gangs' value must be positive" } */
   ;
 #pragma acc parallel num_workers((int) -1.2) /* { dg-warning "'num_workers' value must be positive" } */
@@ -114,7 +218,8 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length((int) -1.2) /* { dg-warning "'vector_length' value must be positive" } */
   ;
 
-#pragma acc parallel \
+
+#pragma acc kernels \
   num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c } } */ \
   num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c } } */ \
   vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c } } */ \
@@ -123,12 +228,31 @@ void acc_parallel(int i, float f)
   num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c++ } } */
   ;
 
-#pragma acc parallel \
+#pragma acc parallel							\
+  num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c } } */ \
+  num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c } } */ \
+  vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c } } */ \
+  num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c++ } } */ \
+  vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c++ } } */ \
+  num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c++ } } */
+  ;
+
+
+#pragma acc kernels \
+  num_gangs(-1) /* { dg-warning "'num_gangs' value must be positive" } */ \
+  num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */ \
+  vector_length(abc_k) /* { dg-error "'abc_k' (un|was not )declared" } */ \
+  num_workers(0.5) /* { dg-error "'num_workers' expression must be integral" } */ \
+  vector_length(&f) /* { dg-error "'vector_length' expression must be integral" } */ \
+  num_gangs( /* { dg-error "expected (primary-|)expression before end of line" "TODO" { xfail c } } */
+  ;
+
+#pragma acc parallel							\
   num_gangs(-1) /* { dg-warning "'num_gangs' value must be positive" } */ \
   num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */ \
-  vector_length(abc) /* { dg-error "'abc' (un|was not )declared" } */ \
+  vector_length(abc_p) /* { dg-error "'abc_p' (un|was not )declared" } */ \
   num_workers(0.5) /* { dg-error "'num_workers' expression must be integral" } */ \
-  vector_length(&acc_parallel) /* { dg-error "'vector_length' expression must be integral" } */ \
+  vector_length(&f) /* { dg-error "'vector_length' expression must be integral" } */ \
   num_gangs( /* { dg-error "expected (primary-|)expression before end of line" "TODO" { xfail c } } */
   ;
 }
diff --git gcc/testsuite/c-c++-common/goacc/routine-1.c gcc/testsuite/c-c++-common/goacc/routine-1.c
index a5e0d69..a756922 100644
--- gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -21,6 +21,13 @@ void seq (void)
 
 int main ()
 {
+#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32)
+  {
+    gang ();
+    worker ();
+    vector ();
+    seq ();
+  }
 
 #pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
   {
diff --git gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
index 0a006e3..9f11196 100644
--- gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
+++ gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
@@ -1,10 +1,6 @@
-/* { dg-do compile } */
 /* { dg-additional-options "-Wuninitialized" } */
 
-#include <stdbool.h>
-
-int
-main (void)
+void acc_parallel()
 {
   int i, j, k;
 
@@ -17,3 +13,17 @@ main (void)
   #pragma acc parallel vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
   ;
 }
+
+void acc_kernels()
+{
+  int i, j, k;
+
+  #pragma acc kernels num_gangs(i) /* { dg-warning "is used uninitialized in this function" } */
+  ;
+
+  #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized in this function" } */
+  ;
+
+  #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
+  ;
+}
diff --git gcc/testsuite/g++.dg/goacc/template.C gcc/testsuite/g++.dg/goacc/template.C
index 74f40d8..852f42f 100644
--- gcc/testsuite/g++.dg/goacc/template.C
+++ gcc/testsuite/g++.dg/goacc/template.C
@@ -100,6 +100,10 @@ oacc_kernels_copy (T a)
   float y = 3;
   double z = 4;
 
+#pragma acc kernels num_gangs (a) num_workers (a) vector_length (a) default (none) copyout (b) copyin (a)
+  for (int i = 0; i < 1; i++)
+    b = a;
+
 #pragma acc kernels copy (w, x, y, z)
   {
     w = accDouble<char>(w);
diff --git gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
index 4ec66de..7daca59 100644
--- gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -6,7 +6,8 @@ program test
   integer :: q, i, j, k, m, n, o, p, r, s, t, u, v, w
   logical :: l = .true.
 
-  !$acc kernels if(l) async copy(i), copyin(j), copyout(k), create(m) &
+  !$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
+  !$acc copy(i), copyin(j), copyout(k), create(m) &
   !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
   !$acc deviceptr(u)
   !$acc end kernels
@@ -16,6 +17,9 @@ end program test
 
 ! { dg-final { scan-tree-dump-times "if" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "async" 1 "original" } } 
+! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } 
+! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } 
+! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } 
 
 ! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } 
 ! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } 
diff --git gcc/testsuite/gfortran.dg/goacc/routine-3.f90 gcc/testsuite/gfortran.dg/goacc/routine-3.f90
index ca9b928..6773f62 100644
--- gcc/testsuite/gfortran.dg/goacc/routine-3.f90
+++ gcc/testsuite/gfortran.dg/goacc/routine-3.f90
@@ -4,6 +4,12 @@ CONTAINS
     INTEGER  :: i
     REAL(KIND=8), ALLOCATABLE :: un(:),  ua(:)
 
+    !$acc kernels num_gangs(2) num_workers(4) vector_length(32)
+    DO jj = 1, 100
+       un(i) = ua(i)
+    END DO
+    !$acc end kernels
+
     !$acc parallel num_gangs(2) num_workers(4) vector_length(32)
     DO jj = 1, 100
        un(i) = ua(i)
diff --git gcc/testsuite/gfortran.dg/goacc/sie.f95 gcc/testsuite/gfortran.dg/goacc/sie.f95
index 2d66026..abfe28b 100644
--- gcc/testsuite/gfortran.dg/goacc/sie.f95
+++ gcc/testsuite/gfortran.dg/goacc/sie.f95
@@ -95,6 +95,34 @@ program test
   !$acc parallel num_gangs("1") ! { dg-error "scalar INTEGER expression" }
   !$acc end parallel
 
+  !$acc kernels num_gangs ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc kernels num_gangs(3)
+  !$acc end kernels
+
+  !$acc kernels num_gangs(i)
+  !$acc end kernels
+
+  !$acc kernels num_gangs(i+1)
+  !$acc end kernels
+
+  !$acc kernels num_gangs(-1) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs(0) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs() ! { dg-error "Invalid character in name" }
+
+  !$acc kernels num_gangs(1.5) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs(.true.) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs("1") ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
 
   !$acc parallel num_workers ! { dg-error "Unclassifiable OpenACC directive" }
 
@@ -124,6 +152,34 @@ program test
   !$acc parallel num_workers("1") ! { dg-error "scalar INTEGER expression" }
   !$acc end parallel
 
+  !$acc kernels num_workers ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc kernels num_workers(3)
+  !$acc end kernels
+
+  !$acc kernels num_workers(i)
+  !$acc end kernels
+
+  !$acc kernels num_workers(i+1)
+  !$acc end kernels
+
+  !$acc kernels num_workers(-1) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_workers(0) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_workers() ! { dg-error "Invalid character in name" }
+
+  !$acc kernels num_workers(1.5) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_workers(.true.) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_workers("1") ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
 
   !$acc parallel vector_length ! { dg-error "Unclassifiable OpenACC directive" }
 
@@ -153,6 +209,34 @@ program test
   !$acc parallel vector_length("1") ! { dg-error "scalar INTEGER expression" }
   !$acc end parallel
 
+  !$acc kernels vector_length ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc kernels vector_length(3)
+  !$acc end kernels
+
+  !$acc kernels vector_length(i)
+  !$acc end kernels
+
+  !$acc kernels vector_length(i+1)
+  !$acc end kernels
+
+  !$acc kernels vector_length(-1) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels vector_length(0) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels vector_length() ! { dg-error "Invalid character in name" }
+
+  !$acc kernels vector_length(1.5) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels vector_length(.true.) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels vector_length("1") ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
 
   !$acc loop gang
   do i = 1,10
@@ -249,4 +333,4 @@ program test
   do i = 1,10
   enddo
 
-end program test
\ No newline at end of file
+end program test
diff --git gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
index b87d26f..5dea42b 100644
--- gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
+++ gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
@@ -1,7 +1,6 @@
-! { dg-do compile }
 ! { dg-additional-options "-Wuninitialized" }
 
-program test
+subroutine acc_parallel
   implicit none
   integer :: i, j, k
 
@@ -13,5 +12,18 @@ program test
 
   !$acc parallel vector_length(k) ! { dg-warning "is used uninitialized in this function" }
   !$acc end parallel
+end subroutine acc_parallel
 
-end program test
+subroutine acc_kernels
+  implicit none
+  integer :: i, j, k
+
+  !$acc kernels num_gangs(i) ! { dg-warning "is used uninitialized in this function" }
+  !$acc end kernels
+
+  !$acc kernels num_workers(j) ! { dg-warning "is used uninitialized in this function" }
+  !$acc end kernels
+
+  !$acc kernels vector_length(k) ! { dg-warning "is used uninitialized in this function" }
+  !$acc end kernels
+end subroutine acc_kernels
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
index c7592d6..b840888 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
@@ -14,27 +14,40 @@ main (void)
   b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
   c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
 
+  /* Parallelism dimensions: compiler/runtime decides.  */
 #pragma acc kernels copyout (a[0:N])
   {
     for (COUNTERTYPE i = 0; i < N; i++)
       a[i] = i * 2;
   }
 
-#pragma acc kernels copyout (b[0:N])
+  /* Parallelism dimensions: variable.  */
+#pragma acc kernels copyout (b[0:N]) \
+  num_gangs (3 + a[3]) num_workers (5 + a[5]) vector_length (7 + a[7])
+  /* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */
   {
     for (COUNTERTYPE i = 0; i < N; i++)
       b[i] = i * 4;
   }
 
-#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
+  /* Parallelism dimensions: literal.  */
+#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) \
+  num_gangs (3) num_workers (5) vector_length (7)
+  /* { dg-prune-output "using vector_length \\(32\\), ignoring 7" } */
   {
     for (COUNTERTYPE ii = 0; ii < N; ii++)
       c[ii] = a[ii] + b[ii];
   }
 
   for (COUNTERTYPE i = 0; i < N; i++)
-    if (c[i] != a[i] + b[i])
-      abort ();
+    {
+      if (a[i] != i * 2)
+	abort ();
+      if (b[i] != i * 4)
+	abort ();
+      if (c[i] != a[i] + b[i])
+	abort ();
+    }
 
   free (a);
   free (b);
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
index d8af546..8308f7c 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
@@ -520,5 +520,40 @@ int main ()
   }
 
 
+  /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1
+     kernels even when there are explicit num_gangs, num_workers, or
+     vector_length clauses.  */
+  {
+    int gangs = 5;
+#define WORKERS 5
+#define VECTORS 13
+    int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
+    gangs_min = workers_min = vectors_min = INT_MAX;
+    gangs_max = workers_max = vectors_max = INT_MIN;
+#pragma acc kernels \
+  num_gangs (gangs) \
+  num_workers (WORKERS) \
+  vector_length (VECTORS)
+    {
+      /* This is to make the OpenACC kernels construct unparallelizable.  */
+      asm volatile ("" : : : "memory");
+
+#pragma acc loop reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
+      for (int i = 100; i > -100; --i)
+	{
+	  gangs_min = gangs_max = acc_gang ();
+	  workers_min = workers_max = acc_worker ();
+	  vectors_min = vectors_max = acc_vector ();
+	}
+    }
+    if (gangs_min != 0 || gangs_max != 1 - 1
+	|| workers_min != 0 || workers_max != 1 - 1
+	|| vectors_min != 0 || vectors_max != 1 - 1)
+      __builtin_abort ();
+#undef VECTORS
+#undef WORKERS
+  }
+
+
   return 0;
 }
diff --git libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95 libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
index 163e8d5..b88ca67 100644
--- libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
+++ libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
@@ -6,25 +6,34 @@ program main
   integer, dimension (0:n-1) :: a, b, c
   integer                    :: i, ii
 
+  ! Parallelism dimensions: compiler/runtime decides.
   !$acc kernels copyout (a(0:n-1))
   do i = 0, n - 1
      a(i) = i * 2
   end do
   !$acc end kernels
 
-  !$acc kernels copyout (b(0:n-1))
+  ! Parallelism dimensions: variable.
+  !$acc kernels copyout (b(0:n-1)) &
+  !$acc num_gangs (3 + a(3)) num_workers (5 + a(5)) vector_length (7 + a(7))
+  ! { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" }
   do i = 0, n -1
      b(i) = i * 4
   end do
   !$acc end kernels
 
-  !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
+  ! Parallelism dimensions: literal.
+  !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) &
+  !$acc num_gangs (3) num_workers (5) vector_length (7)
+  ! { dg-prune-output "using vector_length \\(32\\), ignoring 7" }
   do ii = 0, n - 1
      c(ii) = a(ii) + b(ii)
   end do
   !$acc end kernels
 
   do i = 0, n - 1
+     if (a(i) .ne. i * 2) call abort
+     if (b(i) .ne. i * 4) call abort
      if (c(i) .ne. a(i) + b(i)) call abort
   end do
 


Grüße
 Thomas

^ permalink raw reply	[flat|nested] 5+ messages in thread

* Re: OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses
  2017-05-11 12:28 OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses Thomas Schwinge
@ 2017-05-14 10:40 ` Thomas Schwinge
  2017-05-19 11:09 ` Thomas Schwinge
  2017-05-23  9:54 ` Jakub Jelinek
  2 siblings, 0 replies; 5+ messages in thread
From: Thomas Schwinge @ 2017-05-14 10:40 UTC (permalink / raw)
  To: gcc-patches; +Cc: Jakub Jelinek

Hi!

On Thu, 11 May 2017 14:26:51 +0200, I wrote:
> Building on the other pending patches (I'll soon commit the approved
> ones), we can then support the num_gangs, num_workers, vector_length
> clauses for the OpenACC 2.5 kernels construct.  OK for trunk?

>     OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses

For now, committed to gomp-4_0-branch in r248031:

commit cc2a61ba48e84268e37c53874cb3eef27f5ede1d
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Sun May 14 10:26:07 2017 +0000

    OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses
    
            gcc/c/
            * c-parser.c (OACC_KERNELS_CLAUSE_MASK)
            (OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK): Add
            "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
            "VECTOR_LENGTH".
            gcc/cp/
            * parser.c (OACC_KERNELS_CLAUSE_MASK)
            (OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK): Add
            "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
            "VECTOR_LENGTH".
            gcc/fortran/
            * openmp.c (OACC_KERNELS_CLAUSES)
            (OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK): Add
            "OMP_CLAUSE_NUM_GANGS", "OMP_CLAUSE_NUM_WORKERS",
            "OMP_CLAUSE_VECTOR_LENGTH".
            gcc/
            * omp-low.c (execute_oacc_device_lower): Remove the parallelism
            dimensions function attributes for unparallelized OpenACC kernels
            constructs.
            gcc/testsuite/
            * c-c++-common/goacc/parallel-dims-1.c: Update.
            * c-c++-common/goacc/parallel-dims-2.c: Likewise.
            * c-c++-common/goacc/routine-1.c: Likewise.
            * c-c++-common/goacc/uninit-dim-clause.c: Likewise.
            * g++.dg/goacc/template.C: Likewise.
            * gfortran.dg/goacc/kernels-tree.f95: Likewise.
            * gfortran.dg/goacc/routine-3.f90: Likewise.
            * gfortran.dg/goacc/sie.f95: Likewise.
            * gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
            libgomp/
            * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: New
            file.
            * testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: Update.
            * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
            * testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@248031 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp                                 |   6 +
 gcc/c/ChangeLog.gomp                               |   7 +
 gcc/c/c-parser.c                                   |   6 +
 gcc/cp/ChangeLog.gomp                              |   7 +
 gcc/cp/parser.c                                    |   6 +
 gcc/fortran/ChangeLog.gomp                         |   7 +
 gcc/fortran/openmp.c                               |   6 +-
 gcc/omp-low.c                                      |   9 +
 gcc/testsuite/ChangeLog.gomp                       |  12 +
 gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c |   4 +
 gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c | 152 +++++++++++--
 gcc/testsuite/c-c++-common/goacc/routine-1.c       |  13 ++
 .../c-c++-common/goacc/uninit-dim-clause.c         |  17 +-
 gcc/testsuite/g++.dg/goacc/template.C              |   4 +
 gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95   |   6 +-
 gcc/testsuite/gfortran.dg/goacc/routine-3.f90      |   6 +
 gcc/testsuite/gfortran.dg/goacc/sie.f95            |  84 +++++++
 .../gfortran.dg/goacc/uninit-dim-clause.f95        |  18 +-
 libgomp/ChangeLog.gomp                             |   6 +
 .../libgomp.oacc-c-c++-common/acc_prof-kernels-1.c | 244 +++++++++++++++++++++
 .../libgomp.oacc-c-c++-common/kernels-loop-2.c     |  21 +-
 .../libgomp.oacc-c-c++-common/parallel-dims.c      |  35 +++
 .../libgomp.oacc-fortran/kernels-loop-2.f95        |  13 +-
 23 files changed, 661 insertions(+), 28 deletions(-)

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index a754647..a4720c3 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,9 @@
+2017-05-14  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* omp-low.c (execute_oacc_device_lower): Remove the parallelism
+	dimensions function attributes for unparallelized OpenACC kernels
+	constructs.
+
 2017-05-12  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* config/nvptx/nvptx.c (nvptx_goacc_reduction_init): Don't update
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index 3efcc8b..baedcf8 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,10 @@
+2017-05-14  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-parser.c (OACC_KERNELS_CLAUSE_MASK)
+	(OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK): Add
+	"PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
+	"VECTOR_LENGTH".
+
 2017-05-12  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* c-parser.c (c_parser_omp_clause_num_gangs)
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index ef61c5f..afc467d 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -13966,11 +13966,17 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
 #define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK				\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
 #define OACC_PARALLEL_CLAUSE_MASK					\
diff --git gcc/cp/ChangeLog.gomp gcc/cp/ChangeLog.gomp
index 470f4e7..d59e856 100644
--- gcc/cp/ChangeLog.gomp
+++ gcc/cp/ChangeLog.gomp
@@ -1,3 +1,10 @@
+2017-05-14  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* parser.c (OACC_KERNELS_CLAUSE_MASK)
+	(OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK): Add
+	"PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
+	"VECTOR_LENGTH".
+
 2017-05-04  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* parser.c (cp_parser_omp_clause_name): Add support for if_present.
diff --git gcc/cp/parser.c gcc/cp/parser.c
index b9c9747..de42cdd 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -35704,11 +35704,17 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
 #define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK				\
 	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT))
 
 #define OACC_PARALLEL_CLAUSE_MASK					\
diff --git gcc/fortran/ChangeLog.gomp gcc/fortran/ChangeLog.gomp
index e9d8928..8a6ae6a 100644
--- gcc/fortran/ChangeLog.gomp
+++ gcc/fortran/ChangeLog.gomp
@@ -1,3 +1,10 @@
+2017-05-14  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* openmp.c (OACC_KERNELS_CLAUSES)
+	(OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK): Add
+	"OMP_CLAUSE_NUM_GANGS", "OMP_CLAUSE_NUM_WORKERS",
+	"OMP_CLAUSE_VECTOR_LENGTH".
+
 2017-05-04  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* gfortran.h (gfc_omp_clauses): Add if_present member.
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index 12b2430..c7e78bb 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -2049,7 +2049,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
    | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT        \
    | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
 #define OACC_KERNELS_CLAUSES \
-  (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_DEVICEPTR	      \
+  (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS	      \
+   | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
    | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT                 \
    | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT              \
    | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
@@ -2093,7 +2094,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
    | OMP_CLAUSE_VECTOR | OMP_CLAUSE_AUTO | OMP_CLAUSE_SEQ | OMP_CLAUSE_TILE   \
    | OMP_CLAUSE_DEVICE_TYPE)
 #define OACC_KERNELS_CLAUSE_DEVICE_TYPE_MASK \
-  (omp_mask (OMP_CLAUSE_ASYNC) | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
+  (omp_mask (OMP_CLAUSE_ASYNC) | OMP_CLAUSE_NUM_GANGS | OMP_CLAUSE_NUM_WORKERS \
+   | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
 #define OACC_PARALLEL_CLAUSE_DEVICE_TYPE_MASK				      \
   (omp_mask (OMP_CLAUSE_ASYNC) | OMP_CLAUSE_NUM_GANGS | OMP_CLAUSE_NUM_WORKERS \
    | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_WAIT | OMP_CLAUSE_DEVICE_TYPE)
diff --git gcc/omp-low.c gcc/omp-low.c
index 0fbc3ff..ae8b6d9 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -21178,6 +21178,15 @@ execute_oacc_device_lower ()
     = (lookup_attribute ("oacc kernels parallelized",
 			 DECL_ATTRIBUTES (current_function_decl)) != NULL);
 
+  /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1
+     kernels, so remove the parallelism dimensions function attributes
+     potentially set earlier on.  */
+  if (is_oacc_kernels && !is_oacc_kernels_parallelized)
+    {
+      set_oacc_fn_attrib (current_function_decl, NULL, NULL);
+      attrs = get_oacc_fn_attrib (current_function_decl);
+    }
+
   /* Discover, partition and process the loops.  */
   oacc_loop *loops = oacc_loop_discovery ();
   int fn_level = oacc_fn_attrib_level (attrs);
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index c24820d..67f01e8 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,15 @@
+2017-05-14  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/parallel-dims-1.c: Update.
+	* c-c++-common/goacc/parallel-dims-2.c: Likewise.
+	* c-c++-common/goacc/routine-1.c: Likewise.
+	* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
+	* g++.dg/goacc/template.C: Likewise.
+	* gfortran.dg/goacc/kernels-tree.f95: Likewise.
+	* gfortran.dg/goacc/routine-3.f90: Likewise.
+	* gfortran.dg/goacc/sie.f95: Likewise.
+	* gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
+
 2017-05-12  Thomas Schwinge  <thomas@codesourcery.com>
 
 	* c-c++-common/goacc/parallel-dims-1.c: New file.
diff --git gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
index 9e4cfaa..6cdbebe 100644
--- gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
+++ gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
@@ -3,6 +3,10 @@
 
 void f(int i)
 {
+#pragma acc kernels \
+  num_gangs(i) num_workers(i) vector_length(i)
+  ;
+
 #pragma acc parallel /* { dg-bogus "region is (gang|worker|vector) partitioned" "" { xfail *-*-* } } */ \
   num_gangs(i) num_workers(i) vector_length(i)
   ;
diff --git gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
index 30a3d17..acfbe7f 100644
--- gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
+++ gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
@@ -1,18 +1,15 @@
 /* Invalid use of OpenACC parallelism dimensions clauses: num_gangs,
    num_workers, vector_length.  */
 
-void acc_kernels(int i)
+void f(int i, float f)
 {
-#pragma acc kernels num_gangs(i) /* { dg-error "'num_gangs' is not valid for '#pragma acc kernels'" } */
+#pragma acc kernels num_gangs /* { dg-error "expected '\\(' before end of line" } */
   ;
-#pragma acc kernels num_workers(i) /* { dg-error "'num_workers' is not valid for '#pragma acc kernels'" } */
+#pragma acc kernels num_workers /* { dg-error "expected '\\(' before end of line" } */
   ;
-#pragma acc kernels vector_length(i) /* { dg-error "'vector_length' is not valid for '#pragma acc kernels'" } */
+#pragma acc kernels vector_length /* { dg-error "expected '\\(' before end of line" } */
   ;
-}
 
-void acc_parallel(int i, float f)
-{
 #pragma acc parallel num_gangs /* { dg-error "expected '\\(' before end of line" } */
   ;
 #pragma acc parallel num_workers /* { dg-error "expected '\\(' before end of line" } */
@@ -20,6 +17,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length /* { dg-error "expected '\\(' before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs( /* { dg-error "expected (primary-|)expression before end of line" } */
+  ;
+#pragma acc kernels num_workers( /* { dg-error "expected (primary-|)expression before end of line" } */
+  ;
+#pragma acc kernels vector_length( /* { dg-error "expected (primary-|)expression before end of line" } */
+  ;
+
 #pragma acc parallel num_gangs( /* { dg-error "expected (primary-|)expression before end of line" } */
   ;
 #pragma acc parallel num_workers( /* { dg-error "expected (primary-|)expression before end of line" } */
@@ -27,6 +32,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length( /* { dg-error "expected (primary-|)expression before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
+  ;
+#pragma acc kernels num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
+  ;
+#pragma acc kernels vector_length() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
+  ;
+
 #pragma acc parallel num_gangs() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
   ;
 #pragma acc parallel num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
@@ -34,6 +47,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
   ;
 
+
+#pragma acc kernels num_gangs(1 /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels num_workers(1 /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels vector_length(1 /* { dg-error "expected '\\)' before end of line" } */
+  ;
+
 #pragma acc parallel num_gangs(1 /* { dg-error "expected '\\)' before end of line" } */
   ;
 #pragma acc parallel num_workers(1 /* { dg-error "expected '\\)' before end of line" } */
@@ -41,6 +62,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1 /* { dg-error "expected '\\)' before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs(i /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels num_workers(i /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels vector_length(i /* { dg-error "expected '\\)' before end of line" } */
+  ;
+
 #pragma acc parallel num_gangs(i /* { dg-error "expected '\\)' before end of line" } */
   ;
 #pragma acc parallel num_workers(i /* { dg-error "expected '\\)' before end of line" } */
@@ -48,6 +77,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(i /* { dg-error "expected '\\)' before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs(1 i /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels num_workers(1 i /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels vector_length(1 i /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+
 #pragma acc parallel num_gangs(1 i /* { dg-error "expected '\\)' before 'i'" } */
   ;
 #pragma acc parallel num_workers(1 i /* { dg-error "expected '\\)' before 'i'" } */
@@ -55,6 +92,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1 i /* { dg-error "expected '\\)' before 'i'" } */
   ;
 
+
+#pragma acc kernels num_gangs(1 i) /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels num_workers(1 i) /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels vector_length(1 i) /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+
 #pragma acc parallel num_gangs(1 i) /* { dg-error "expected '\\)' before 'i'" } */
   ;
 #pragma acc parallel num_workers(1 i) /* { dg-error "expected '\\)' before 'i'" } */
@@ -62,6 +107,17 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1 i) /* { dg-error "expected '\\)' before 'i'" } */
   ;
 
+
+#pragma acc kernels num_gangs(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
+  ;
+#pragma acc kernels num_workers(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
+  ;
+#pragma acc kernels vector_length(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
+  ;
+
 #pragma acc parallel num_gangs(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
   /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
   ;
@@ -72,6 +128,14 @@ void acc_parallel(int i, float f)
   /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
   ;
 
+
+#pragma acc kernels num_gangs(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  ;
+#pragma acc kernels num_workers(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  ;
+#pragma acc kernels vector_length(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  ;
+
 #pragma acc parallel num_gangs(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
   ;
 #pragma acc parallel num_workers(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
@@ -79,11 +143,27 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
   ;
 
-#pragma acc parallel num_gangs(num_gangs) /* { dg-error "'num_gangs' (un|was not )declared" } */
+
+#pragma acc kernels num_gangs(num_gangs_k) /* { dg-error "'num_gangs_k' (un|was not )declared" } */
   ;
-#pragma acc parallel num_workers(num_workers) /* { dg-error "'num_workers' (un|was not )declared" } */
+#pragma acc kernels num_workers(num_workers_k) /* { dg-error "'num_workers_k' (un|was not )declared" } */
   ;
-#pragma acc parallel vector_length(vector_length) /* { dg-error "'vector_length' (un|was not )declared" } */
+#pragma acc kernels vector_length(vector_length_k) /* { dg-error "'vector_length_k' (un|was not )declared" } */
+  ;
+
+#pragma acc parallel num_gangs(num_gangs_p) /* { dg-error "'num_gangs_p' (un|was not )declared" } */
+  ;
+#pragma acc parallel num_workers(num_workers_p) /* { dg-error "'num_workers_p' (un|was not )declared" } */
+  ;
+#pragma acc parallel vector_length(vector_length_p) /* { dg-error "'vector_length_p' (un|was not )declared" } */
+  ;
+
+
+#pragma acc kernels num_gangs(f) /* { dg-error "'num_gangs' expression must be integral" } */
+  ;
+#pragma acc kernels num_workers(f) /* { dg-error "'num_workers' expression must be integral" } */
+  ;
+#pragma acc kernels vector_length(f) /* { dg-error "'vector_length' expression must be integral" } */
   ;
 
 #pragma acc parallel num_gangs(f) /* { dg-error "'num_gangs' expression must be integral" } */
@@ -93,6 +173,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(f) /* { dg-error "'vector_length' expression must be integral" } */
   ;
 
+
+#pragma acc kernels num_gangs((float) 1) /* { dg-error "'num_gangs' expression must be integral" } */
+  ;
+#pragma acc kernels num_workers((float) 1) /* { dg-error "'num_workers' expression must be integral" } */
+  ;
+#pragma acc kernels vector_length((float) 1) /* { dg-error "'vector_length' expression must be integral" } */
+  ;
+
 #pragma acc parallel num_gangs((float) 1) /* { dg-error "'num_gangs' expression must be integral" } */
   ;
 #pragma acc parallel num_workers((float) 1) /* { dg-error "'num_workers' expression must be integral" } */
@@ -100,6 +188,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length((float) 1) /* { dg-error "'vector_length' expression must be integral" } */
   ;
 
+
+#pragma acc kernels num_gangs(0) /* { dg-warning "'num_gangs' value must be positive" } */
+  ;
+#pragma acc kernels num_workers(0) /* { dg-warning "'num_workers' value must be positive" } */
+  ;
+#pragma acc kernels vector_length(0) /* { dg-warning "'vector_length' value must be positive" } */
+  ;
+
 #pragma acc parallel num_gangs(0) /* { dg-warning "'num_gangs' value must be positive" } */
   ;
 #pragma acc parallel num_workers(0) /* { dg-warning "'num_workers' value must be positive" } */
@@ -107,6 +203,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(0) /* { dg-warning "'vector_length' value must be positive" } */
   ;
 
+
+#pragma acc kernels num_gangs((int) -1.2) /* { dg-warning "'num_gangs' value must be positive" } */
+  ;
+#pragma acc kernels num_workers((int) -1.2) /* { dg-warning "'num_workers' value must be positive" } */
+  ;
+#pragma acc kernels vector_length((int) -1.2) /* { dg-warning "'vector_length' value must be positive" } */
+  ;
+
 #pragma acc parallel num_gangs((int) -1.2) /* { dg-warning "'num_gangs' value must be positive" } */
   ;
 #pragma acc parallel num_workers((int) -1.2) /* { dg-warning "'num_workers' value must be positive" } */
@@ -114,7 +218,8 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length((int) -1.2) /* { dg-warning "'vector_length' value must be positive" } */
   ;
 
-#pragma acc parallel \
+
+#pragma acc kernels \
   num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c } } */ \
   num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c } } */ \
   vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c } } */ \
@@ -123,12 +228,31 @@ void acc_parallel(int i, float f)
   num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c++ } } */
   ;
 
-#pragma acc parallel \
+#pragma acc parallel							\
+  num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c } } */ \
+  num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c } } */ \
+  vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c } } */ \
+  num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c++ } } */ \
+  vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c++ } } */ \
+  num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c++ } } */
+  ;
+
+
+#pragma acc kernels \
+  num_gangs(-1) /* { dg-warning "'num_gangs' value must be positive" } */ \
+  num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */ \
+  vector_length(abc_k) /* { dg-error "'abc_k' (un|was not )declared" } */ \
+  num_workers(0.5) /* { dg-error "'num_workers' expression must be integral" } */ \
+  vector_length(&f) /* { dg-error "'vector_length' expression must be integral" } */ \
+  num_gangs( /* { dg-error "expected (primary-|)expression before end of line" "TODO" { xfail c } } */
+  ;
+
+#pragma acc parallel							\
   num_gangs(-1) /* { dg-warning "'num_gangs' value must be positive" } */ \
   num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */ \
-  vector_length(abc) /* { dg-error "'abc' (un|was not )declared" } */ \
+  vector_length(abc_p) /* { dg-error "'abc_p' (un|was not )declared" } */ \
   num_workers(0.5) /* { dg-error "'num_workers' expression must be integral" } */ \
-  vector_length(&acc_parallel) /* { dg-error "'vector_length' expression must be integral" } */ \
+  vector_length(&f) /* { dg-error "'vector_length' expression must be integral" } */ \
   num_gangs( /* { dg-error "expected (primary-|)expression before end of line" "TODO" { xfail c } } */
   ;
 }
diff --git gcc/testsuite/c-c++-common/goacc/routine-1.c gcc/testsuite/c-c++-common/goacc/routine-1.c
index a4ecfd3..7389575 100644
--- gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -91,6 +91,19 @@ extern void nohost (void);
 
 int main ()
 {
+#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32)
+  {
+    gang ();
+    worker ();
+    vector ();
+    seq ();
+    bind_f_1 ();
+    bind_f_1_1 ();
+    bind_f_2 ();
+    bind_f_2_1 ();
+    nohost ();
+  }
+
 #pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
   {
     gang ();
diff --git gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
index b6e4c3d..72aacd7 100644
--- gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
+++ gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
@@ -1,7 +1,6 @@
 /* { dg-additional-options "-Wuninitialized" } */
 
-int
-main (void)
+void acc_parallel()
 {
   int i, j, k;
 
@@ -16,6 +15,18 @@ main (void)
   #pragma acc parallel loop vector vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
   for (k = 0; k < 1; k++)
     ;
+}
+
+void acc_kernels()
+{
+  int i, j, k;
+
+  #pragma acc kernels num_gangs(i) /* { dg-warning "is used uninitialized in this function" } */
+  ;
+
+  #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized in this function" } */
+  ;
 
-  return 0;
+  #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
+  ;
 }
diff --git gcc/testsuite/g++.dg/goacc/template.C gcc/testsuite/g++.dg/goacc/template.C
index 4bc2596..f4d255c 100644
--- gcc/testsuite/g++.dg/goacc/template.C
+++ gcc/testsuite/g++.dg/goacc/template.C
@@ -100,6 +100,10 @@ oacc_kernels_copy (T a)
   float y = 3;
   double z = 4;
 
+#pragma acc kernels num_gangs (a) num_workers (a) vector_length (a) default (none) copyout (b) copyin (a)
+  for (int i = 0; i < 1; i++)
+    b = a;
+
 #pragma acc kernels copy (w, x, y, z)
   {
     w = accDouble<char>(w);
diff --git gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
index 2c237b7..a70f1e7 100644
--- gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -6,7 +6,8 @@ program test
   integer :: q, i, j, k, m, n, o, p, r, s, t, u, v, w
   logical :: l = .true.
 
-  !$acc kernels if(l) async copy(i), copyin(j), copyout(k), create(m) &
+  !$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
+  !$acc copy(i), copyin(j), copyout(k), create(m) &
   !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
   !$acc deviceptr(u)
   !$acc end kernels
@@ -16,6 +17,9 @@ end program test
 
 ! { dg-final { scan-tree-dump-times "if" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "async" 1 "original" } } 
+! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } 
+! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } 
+! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } 
 
 ! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } 
 ! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } 
diff --git gcc/testsuite/gfortran.dg/goacc/routine-3.f90 gcc/testsuite/gfortran.dg/goacc/routine-3.f90
index ca9b928..6773f62 100644
--- gcc/testsuite/gfortran.dg/goacc/routine-3.f90
+++ gcc/testsuite/gfortran.dg/goacc/routine-3.f90
@@ -4,6 +4,12 @@ CONTAINS
     INTEGER  :: i
     REAL(KIND=8), ALLOCATABLE :: un(:),  ua(:)
 
+    !$acc kernels num_gangs(2) num_workers(4) vector_length(32)
+    DO jj = 1, 100
+       un(i) = ua(i)
+    END DO
+    !$acc end kernels
+
     !$acc parallel num_gangs(2) num_workers(4) vector_length(32)
     DO jj = 1, 100
        un(i) = ua(i)
diff --git gcc/testsuite/gfortran.dg/goacc/sie.f95 gcc/testsuite/gfortran.dg/goacc/sie.f95
index b4dd9ed..3abf2c8 100644
--- gcc/testsuite/gfortran.dg/goacc/sie.f95
+++ gcc/testsuite/gfortran.dg/goacc/sie.f95
@@ -95,6 +95,34 @@ program test
   !$acc parallel num_gangs("1") ! { dg-error "scalar INTEGER expression" }
   !$acc end parallel
 
+  !$acc kernels num_gangs ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc kernels num_gangs(3)
+  !$acc end kernels
+
+  !$acc kernels num_gangs(i)
+  !$acc end kernels
+
+  !$acc kernels num_gangs(i+1)
+  !$acc end kernels
+
+  !$acc kernels num_gangs(-1) ! { dg-error "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs(0) ! { dg-error "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs() ! { dg-error "Invalid character in name" }
+
+  !$acc kernels num_gangs(1.5) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs(.true.) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs("1") ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
 
   !$acc parallel num_workers ! { dg-error "Unclassifiable OpenACC directive" }
 
@@ -124,6 +152,34 @@ program test
   !$acc parallel num_workers("1") ! { dg-error "scalar INTEGER expression" }
   !$acc end parallel
 
+  !$acc kernels num_workers ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc kernels num_workers(3)
+  !$acc end kernels
+
+  !$acc kernels num_workers(i)
+  !$acc end kernels
+
+  !$acc kernels num_workers(i+1)
+  !$acc end kernels
+
+  !$acc kernels num_workers(-1) ! { dg-error "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_workers(0) ! { dg-error "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_workers() ! { dg-error "Invalid character in name" }
+
+  !$acc kernels num_workers(1.5) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_workers(.true.) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_workers("1") ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
 
   !$acc parallel vector_length ! { dg-error "Unclassifiable OpenACC directive" }
 
@@ -153,6 +209,34 @@ program test
   !$acc parallel vector_length("1") ! { dg-error "scalar INTEGER expression" }
   !$acc end parallel
 
+  !$acc kernels vector_length ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc kernels vector_length(3)
+  !$acc end kernels
+
+  !$acc kernels vector_length(i)
+  !$acc end kernels
+
+  !$acc kernels vector_length(i+1)
+  !$acc end kernels
+
+  !$acc kernels vector_length(-1) ! { dg-error "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels vector_length(0) ! { dg-error "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels vector_length() ! { dg-error "Invalid character in name" }
+
+  !$acc kernels vector_length(1.5) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels vector_length(.true.) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels vector_length("1") ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
 
   !$acc loop gang
   do i = 1,10
diff --git gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
index 9db692a..8551140 100644
--- gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
+++ gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
@@ -1,6 +1,6 @@
 ! { dg-additional-options "-Wuninitialized" }
 
-program test
+subroutine acc_parallel
   implicit none
   integer :: i, j, k
 
@@ -18,4 +18,18 @@ program test
   do k = 0, 1
   end do
   !$acc end parallel loop
-end program test
+end subroutine acc_parallel
+
+subroutine acc_kernels
+  implicit none
+  integer :: i, j, k
+
+  !$acc kernels num_gangs(i) ! { dg-warning "is used uninitialized in this function" }
+  !$acc end kernels
+
+  !$acc kernels num_workers(j) ! { dg-warning "is used uninitialized in this function" }
+  !$acc end kernels
+
+  !$acc kernels vector_length(k) ! { dg-warning "is used uninitialized in this function" }
+  !$acc end kernels
+end subroutine acc_kernels
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index a1627a8..5dc0889 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,5 +1,11 @@
 2017-05-14  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: New
+	file.
+	* testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: Update.
+	* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
+	* testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise.
+
 	* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Rewrite.
 	* testsuite/libgomp.oacc-c++/c++.exp (check_effective_target_c)
 	(check_effective_target_c++): New procs.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
new file mode 100644
index 0000000..24b5718
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
@@ -0,0 +1,244 @@
+/* Test dispatch of events to callbacks.  */
+
+#undef NDEBUG
+#include <assert.h>
+#include <stdlib.h>
+#include <string.h>
+
+#include <acc_prof.h>
+
+/* Use explicit copyin clauses, to work around firstprivate optimizations,
+   which will cause the value at the point of call to be used (*before* any
+   potential modifications done in callbacks), as opposed to its address being
+   taken, which then later gets dereferenced (*after* any modifications done in
+   callbacks).  */
+#define COPYIN(...) copyin(__VA_ARGS__)
+
+/* See the "DEVICE_INIT_INSIDE_COMPUTE_CONSTRUCT" reference in
+   libgomp.texi.  */
+#define DEVICE_INIT_INSIDE_COMPUTE_CONSTRUCT 0
+
+#define DEBUG_printf(...) //__builtin_printf (__VA_ARGS__)
+
+static int state = -1;
+#define STATE_OP(state, op)\
+  do \
+    { \
+      typeof (state) state_o = (state); \
+      (void) state_o; \
+      (state)op; \
+      DEBUG_printf("state: %d -> %d\n", state_o, (state)); \
+    } \
+  while (0)
+
+static acc_device_t acc_device_type;
+static int acc_device_num;
+static int num_gangs, num_workers, vector_length;
+
+void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
+{
+  DEBUG_printf ("%s\n", __FUNCTION__);
+
+  assert (acc_device_type != acc_device_host);
+
+  assert (state == 0);
+  STATE_OP (state, = 1);
+
+  assert (prof_info->event_type == acc_ev_enqueue_launch_start);
+  assert (prof_info->valid_bytes == _ACC_PROF_INFO_VALID_BYTES);
+  assert (prof_info->version == _ACC_PROF_INFO_VERSION);
+  assert (prof_info->device_type == acc_device_type);
+  assert (prof_info->device_number == acc_device_num);
+  assert (prof_info->thread_id == -1);
+  assert (prof_info->async == acc_async_sync);
+  assert (prof_info->async_queue == prof_info->async);
+  assert (prof_info->src_file == NULL);
+  assert (prof_info->func_name == NULL);
+  assert (prof_info->line_no == -1);
+  assert (prof_info->end_line_no == -1);
+  assert (prof_info->func_line_no == -1);
+  assert (prof_info->func_end_line_no == -1);
+
+  assert (event_info->launch_event.event_type == prof_info->event_type);
+  assert (event_info->launch_event.valid_bytes == _ACC_LAUNCH_EVENT_INFO_VALID_BYTES);
+  assert (event_info->launch_event.parent_construct == acc_construct_parallel);
+  assert (event_info->launch_event.implicit == 1);
+  assert (event_info->launch_event.tool_info == NULL);
+  assert (event_info->launch_event.kernel_name != NULL);
+  {
+    char *s = strstr (event_info->launch_event.kernel_name, "main");
+    assert (s != NULL);
+    s = strstr (s, "omp_fn");
+    assert (s != NULL);
+  }
+  if (num_gangs < 1)
+    assert (event_info->launch_event.num_gangs >= 1);
+  else
+    {
+#ifdef __OPTIMIZE__
+      assert (event_info->launch_event.num_gangs == num_gangs);
+#else
+      /* No parallelized OpenACC kernels constructs, and unparallelized OpenACC
+	 kernels constructs must get launched as 1 x 1 x 1 kernels.  */
+      assert (event_info->launch_event.num_gangs == 1);
+#endif
+    }
+  if (num_workers < 1)
+    assert (event_info->launch_event.num_workers >= 1);
+  else
+    {
+#ifdef __OPTIMIZE__
+      assert (event_info->launch_event.num_workers == num_workers);
+#else
+      /* See num_gangs above.  */
+      assert (event_info->launch_event.num_workers == 1);
+#endif
+    }
+  if (vector_length < 1)
+    assert (event_info->launch_event.vector_length >= 1);
+  else if (acc_device_type == acc_device_nvidia) /* ... is special.  */
+    assert (event_info->launch_event.vector_length == 32);
+  else
+    {
+#ifdef __OPTIMIZE__
+      assert (event_info->launch_event.vector_length == vector_length);
+#else
+      /* See num_gangs above.  */
+      assert (event_info->launch_event.vector_length == 1);
+#endif
+    }
+
+  if (acc_device_type == acc_device_host)
+    assert (api_info->device_api == acc_device_api_none);
+  else
+    assert (api_info->device_api == acc_device_api_cuda);
+  assert (api_info->valid_bytes == _ACC_API_INFO_VALID_BYTES);
+  assert (api_info->device_type == prof_info->device_type);
+  assert (api_info->vendor == -1);
+  assert (api_info->device_handle == NULL);
+  assert (api_info->context_handle == NULL);
+  assert (api_info->async_handle == NULL);
+}
+
+static acc_prof_reg reg;
+static acc_prof_reg unreg;
+static acc_prof_lookup_func lookup;
+void acc_register_library (acc_prof_reg reg_, acc_prof_reg unreg_, acc_prof_lookup_func lookup_)
+{
+  DEBUG_printf ("%s\n", __FUNCTION__);
+
+  reg = reg_;
+  unreg = unreg_;
+  lookup = lookup_;
+}
+
+int main()
+{
+  STATE_OP (state, = 0);
+  reg (acc_ev_enqueue_launch_start, cb_enqueue_launch_start, acc_reg);
+  assert (state == 0);
+
+  acc_device_type = acc_get_device_type ();
+  acc_device_num = acc_get_device_num (acc_device_type);
+  assert (state == 0);
+
+  /* Parallelism dimensions: compiler/runtime decides.  */
+  STATE_OP (state, = 0);
+  num_gangs = num_workers = vector_length = 0;
+  {
+#define N 100
+    int x[N];
+#pragma acc kernels
+    {
+      for (int i = 0; i < N; ++i)
+	x[i] = i * i;
+    }
+#ifdef __OPTIMIZE__
+    /* TODO.  With -O2 optimizations enabled, the compiler believes that here
+       "state == 0" still holds.  It's not yet clear what's going on.
+       Mis-optimization across the GOMP function call boundary?  Per its
+       gcc/omp-builtins.def definition, BUILT_IN_GOACC_PARALLEL
+       "GOACC_parallel_keyed" doesn't have a "leaf" attribute, so the compiler
+       must expect calls back into this compilation unit?  */
+    asm volatile ("" : : : "memory");
+#endif
+    if (acc_device_type == acc_device_host)
+      assert (state == 0); /* No acc_ev_enqueue_launch_start.  */
+    else
+      assert (state == 1);
+    for (int i = 0; i < N; ++i)
+      if (x[i] != i * i)
+	__builtin_abort ();
+#undef N
+  }
+
+  /* Parallelism dimensions: literal.  */
+  STATE_OP (state, = 0);
+  num_gangs = 30;
+  num_workers = 3;
+  vector_length = 5;
+  {
+#define N 100
+    int x[N];
+#pragma acc kernels \
+  num_gangs (30) num_workers (3) vector_length (5)
+    /* { dg-prune-output "using vector_length \\(32\\), ignoring 5" } */
+    {
+      for (int i = 0; i < N; ++i)
+	x[i] = i * i;
+    }
+#ifdef __OPTIMIZE__
+    /* TODO.  With -O2 optimizations enabled, the compiler believes that here
+       "state == 0" still holds.  It's not yet clear what's going on.
+       Mis-optimization across the GOMP function call boundary?  Per its
+       gcc/omp-builtins.def definition, BUILT_IN_GOACC_PARALLEL
+       "GOACC_parallel_keyed" doesn't have a "leaf" attribute, so the compiler
+       must expect calls back into this compilation unit?  */
+    asm volatile ("" : : : "memory");
+#endif
+    if (acc_device_type == acc_device_host)
+      assert (state == 0); /* No acc_ev_enqueue_launch_start.  */
+    else
+      assert (state == 1);
+    for (int i = 0; i < N; ++i)
+      if (x[i] != i * i)
+	__builtin_abort ();
+#undef N
+  }
+
+  /* Parallelism dimensions: variable.  */
+  STATE_OP (state, = 0);
+  num_gangs = 22;
+  num_workers = 5;
+  vector_length = 7;
+  {
+#define N 100
+    int x[N];
+#pragma acc kernels \
+  num_gangs (num_gangs) num_workers (num_workers) vector_length (vector_length)
+    /* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */
+    {
+      for (int i = 0; i < N; ++i)
+	x[i] = i * i;
+    }
+#ifdef __OPTIMIZE__
+    /* TODO.  With -O2 optimizations enabled, the compiler believes that here
+       "state == 0" still holds.  It's not yet clear what's going on.
+       Mis-optimization across the GOMP function call boundary?  Per its
+       gcc/omp-builtins.def definition, BUILT_IN_GOACC_PARALLEL
+       "GOACC_parallel_keyed" doesn't have a "leaf" attribute, so the compiler
+       must expect calls back into this compilation unit?  */
+    asm volatile ("" : : : "memory");
+#endif
+    if (acc_device_type == acc_device_host)
+      assert (state == 0); /* No acc_ev_enqueue_launch_start.  */
+    else
+      assert (state == 1);
+    for (int i = 0; i < N; ++i)
+      if (x[i] != i * i)
+	__builtin_abort ();
+#undef N
+  }
+
+  return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
index c7592d6..b840888 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
@@ -14,27 +14,40 @@ main (void)
   b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
   c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
 
+  /* Parallelism dimensions: compiler/runtime decides.  */
 #pragma acc kernels copyout (a[0:N])
   {
     for (COUNTERTYPE i = 0; i < N; i++)
       a[i] = i * 2;
   }
 
-#pragma acc kernels copyout (b[0:N])
+  /* Parallelism dimensions: variable.  */
+#pragma acc kernels copyout (b[0:N]) \
+  num_gangs (3 + a[3]) num_workers (5 + a[5]) vector_length (7 + a[7])
+  /* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */
   {
     for (COUNTERTYPE i = 0; i < N; i++)
       b[i] = i * 4;
   }
 
-#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
+  /* Parallelism dimensions: literal.  */
+#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) \
+  num_gangs (3) num_workers (5) vector_length (7)
+  /* { dg-prune-output "using vector_length \\(32\\), ignoring 7" } */
   {
     for (COUNTERTYPE ii = 0; ii < N; ii++)
       c[ii] = a[ii] + b[ii];
   }
 
   for (COUNTERTYPE i = 0; i < N; i++)
-    if (c[i] != a[i] + b[i])
-      abort ();
+    {
+      if (a[i] != i * 2)
+	abort ();
+      if (b[i] != i * 4)
+	abort ();
+      if (c[i] != a[i] + b[i])
+	abort ();
+    }
 
   free (a);
   free (b);
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
index 3458757..1dd6353 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
@@ -523,5 +523,40 @@ int main ()
   }
 
 
+  /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1
+     kernels even when there are explicit num_gangs, num_workers, or
+     vector_length clauses.  */
+  {
+    int gangs = 5;
+#define WORKERS 5
+#define VECTORS 13
+    int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
+    gangs_min = workers_min = vectors_min = INT_MAX;
+    gangs_max = workers_max = vectors_max = INT_MIN;
+#pragma acc kernels \
+  num_gangs (gangs) \
+  num_workers (WORKERS) \
+  vector_length (VECTORS)
+    {
+      /* This is to make the OpenACC kernels construct unparallelizable.  */
+      asm volatile ("" : : : "memory");
+
+#pragma acc loop reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
+      for (int i = 100; i > -100; --i)
+	{
+	  gangs_min = gangs_max = acc_gang ();
+	  workers_min = workers_max = acc_worker ();
+	  vectors_min = vectors_max = acc_vector ();
+	}
+    }
+    if (gangs_min != 0 || gangs_max != 1 - 1
+	|| workers_min != 0 || workers_max != 1 - 1
+	|| vectors_min != 0 || vectors_max != 1 - 1)
+      __builtin_abort ();
+#undef VECTORS
+#undef WORKERS
+  }
+
+
   return 0;
 }
diff --git libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95 libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
index 163e8d5..b88ca67 100644
--- libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
+++ libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
@@ -6,25 +6,34 @@ program main
   integer, dimension (0:n-1) :: a, b, c
   integer                    :: i, ii
 
+  ! Parallelism dimensions: compiler/runtime decides.
   !$acc kernels copyout (a(0:n-1))
   do i = 0, n - 1
      a(i) = i * 2
   end do
   !$acc end kernels
 
-  !$acc kernels copyout (b(0:n-1))
+  ! Parallelism dimensions: variable.
+  !$acc kernels copyout (b(0:n-1)) &
+  !$acc num_gangs (3 + a(3)) num_workers (5 + a(5)) vector_length (7 + a(7))
+  ! { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" }
   do i = 0, n -1
      b(i) = i * 4
   end do
   !$acc end kernels
 
-  !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
+  ! Parallelism dimensions: literal.
+  !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) &
+  !$acc num_gangs (3) num_workers (5) vector_length (7)
+  ! { dg-prune-output "using vector_length \\(32\\), ignoring 7" }
   do ii = 0, n - 1
      c(ii) = a(ii) + b(ii)
   end do
   !$acc end kernels
 
   do i = 0, n - 1
+     if (a(i) .ne. i * 2) call abort
+     if (b(i) .ne. i * 4) call abort
      if (c(i) .ne. a(i) + b(i)) call abort
   end do
 


Grüße
 Thomas

^ permalink raw reply	[flat|nested] 5+ messages in thread

* Re: OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses
  2017-05-11 12:28 OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses Thomas Schwinge
  2017-05-14 10:40 ` Thomas Schwinge
@ 2017-05-19 11:09 ` Thomas Schwinge
  2017-05-23  9:54 ` Jakub Jelinek
  2 siblings, 0 replies; 5+ messages in thread
From: Thomas Schwinge @ 2017-05-19 11:09 UTC (permalink / raw)
  To: gcc-patches, Jakub Jelinek

Hi!

Ping.

On Thu, 11 May 2017 14:26:51 +0200, I wrote:
> [...] support the num_gangs, num_workers, vector_length
> clauses for the OpenACC 2.5 kernels construct.  OK for trunk?
> 
> commit a689c52cde71960bc08ae30c3f88980f66fdd0b8
> Author: Thomas Schwinge <thomas@codesourcery.com>
> Date:   Thu May 11 13:43:28 2017 +0200
> 
>     OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses
>     
>             gcc/c/
>             * c-parser.c (OACC_KERNELS_CLAUSE_MASK): Add
>             "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
>             "VECTOR_LENGTH".
>             gcc/cp/
>             * parser.c (OACC_KERNELS_CLAUSE_MASK): Add
>             "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
>             "VECTOR_LENGTH".
>             gcc/fortran/
>             * openmp.c (OACC_KERNELS_CLAUSES): Add "OMP_CLAUSE_NUM_GANGS",
>             "OMP_CLAUSE_NUM_WORKERS", "OMP_CLAUSE_VECTOR_LENGTH".
>             gcc/
>             * omp-offload.c (execute_oacc_device_lower): Remove the
>             parallelism dimensions function attributes for unparallelized
>             OpenACC kernels constructs.
>             gcc/testsuite/
>             * c-c++-common/goacc/parallel-dims-1.c: Update.
>             * c-c++-common/goacc/parallel-dims-2.c: Likewise.
>             * c-c++-common/goacc/routine-1.c: Likewise.
>             * c-c++-common/goacc/uninit-dim-clause.c: Likewise.
>             * g++.dg/goacc/template.C: Likewise.
>             * gfortran.dg/goacc/kernels-tree.f95: Likewise.
>             * gfortran.dg/goacc/routine-3.f90: Likewise.
>             * gfortran.dg/goacc/sie.f95: Likewise.
>             * gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
>             libgomp/
>             * testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: Update.
>             * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
>             * testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise.
> ---
>  gcc/c/c-parser.c                                   |   3 +
>  gcc/cp/parser.c                                    |   3 +
>  gcc/fortran/openmp.c                               |   3 +-
>  gcc/omp-offload.c                                  |   9 ++
>  gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c |   3 +
>  gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c | 152 +++++++++++++++++++--
>  gcc/testsuite/c-c++-common/goacc/routine-1.c       |   7 +
>  .../c-c++-common/goacc/uninit-dim-clause.c         |  20 ++-
>  gcc/testsuite/g++.dg/goacc/template.C              |   4 +
>  gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95   |   6 +-
>  gcc/testsuite/gfortran.dg/goacc/routine-3.f90      |   6 +
>  gcc/testsuite/gfortran.dg/goacc/sie.f95            |  86 +++++++++++-
>  .../gfortran.dg/goacc/uninit-dim-clause.f95        |  18 ++-
>  .../libgomp.oacc-c-c++-common/kernels-loop-2.c     |  21 ++-
>  .../libgomp.oacc-c-c++-common/parallel-dims.c      |  35 +++++
>  .../libgomp.oacc-fortran/kernels-loop-2.f95        |  13 +-
>  16 files changed, 358 insertions(+), 31 deletions(-)
> 
> diff --git gcc/c/c-parser.c gcc/c/c-parser.c
> index 90d2d17..c0d733c 100644
> --- gcc/c/c-parser.c
> +++ gcc/c/c-parser.c
> @@ -13978,11 +13978,14 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT)		\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
> +	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
> +	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY)	\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)	\
> +	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
>  
>  #define OACC_PARALLEL_CLAUSE_MASK					\
> diff --git gcc/cp/parser.c gcc/cp/parser.c
> index 17d2679..0578e81 100644
> --- gcc/cp/parser.c
> +++ gcc/cp/parser.c
> @@ -36438,11 +36438,14 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT)		\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
> +	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
> +	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY)	\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)	\
> +	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
>  	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
>  
>  #define OACC_PARALLEL_CLAUSE_MASK					\
> diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
> index 89eecfa..7b18a1d 100644
> --- gcc/fortran/openmp.c
> +++ gcc/fortran/openmp.c
> @@ -1926,7 +1926,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
>     | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE \
>     | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
>  #define OACC_KERNELS_CLAUSES \
> -  (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_DEVICEPTR	      \
> +  (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS	      \
> +   | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
>     | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT		      \
>     | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY      \
>     | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT	      \
> diff --git gcc/omp-offload.c gcc/omp-offload.c
> index e954ee9..fab26d0 100644
> --- gcc/omp-offload.c
> +++ gcc/omp-offload.c
> @@ -1452,6 +1452,15 @@ execute_oacc_device_lower ()
>      = (lookup_attribute ("oacc kernels parallelized",
>  			 DECL_ATTRIBUTES (current_function_decl)) != NULL);
>  
> +  /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1
> +     kernels, so remove the parallelism dimensions function attributes
> +     potentially set earlier on.  */
> +  if (is_oacc_kernels && !is_oacc_kernels_parallelized)
> +    {
> +      oacc_set_fn_attrib (current_function_decl, NULL, NULL);
> +      attrs = oacc_get_fn_attrib (current_function_decl);
> +    }
> +
>    /* Discover, partition and process the loops.  */
>    oacc_loop *loops = oacc_loop_discovery ();
>    int fn_level = oacc_fn_attrib_level (attrs);
> diff --git gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
> index a85d3d3..57f682f 100644
> --- gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
> +++ gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
> @@ -3,6 +3,9 @@
>  
>  void f(int i)
>  {
> +#pragma acc kernels num_gangs(i) num_workers(i) vector_length(i)
> +  ;
> +
>  #pragma acc parallel num_gangs(i) num_workers(i) vector_length(i)
>    ;
>  }
> diff --git gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
> index 30a3d17..acfbe7f 100644
> --- gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
> +++ gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
> @@ -1,18 +1,15 @@
>  /* Invalid use of OpenACC parallelism dimensions clauses: num_gangs,
>     num_workers, vector_length.  */
>  
> -void acc_kernels(int i)
> +void f(int i, float f)
>  {
> -#pragma acc kernels num_gangs(i) /* { dg-error "'num_gangs' is not valid for '#pragma acc kernels'" } */
> +#pragma acc kernels num_gangs /* { dg-error "expected '\\(' before end of line" } */
>    ;
> -#pragma acc kernels num_workers(i) /* { dg-error "'num_workers' is not valid for '#pragma acc kernels'" } */
> +#pragma acc kernels num_workers /* { dg-error "expected '\\(' before end of line" } */
>    ;
> -#pragma acc kernels vector_length(i) /* { dg-error "'vector_length' is not valid for '#pragma acc kernels'" } */
> +#pragma acc kernels vector_length /* { dg-error "expected '\\(' before end of line" } */
>    ;
> -}
>  
> -void acc_parallel(int i, float f)
> -{
>  #pragma acc parallel num_gangs /* { dg-error "expected '\\(' before end of line" } */
>    ;
>  #pragma acc parallel num_workers /* { dg-error "expected '\\(' before end of line" } */
> @@ -20,6 +17,14 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length /* { dg-error "expected '\\(' before end of line" } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs( /* { dg-error "expected (primary-|)expression before end of line" } */
> +  ;
> +#pragma acc kernels num_workers( /* { dg-error "expected (primary-|)expression before end of line" } */
> +  ;
> +#pragma acc kernels vector_length( /* { dg-error "expected (primary-|)expression before end of line" } */
> +  ;
> +
>  #pragma acc parallel num_gangs( /* { dg-error "expected (primary-|)expression before end of line" } */
>    ;
>  #pragma acc parallel num_workers( /* { dg-error "expected (primary-|)expression before end of line" } */
> @@ -27,6 +32,14 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length( /* { dg-error "expected (primary-|)expression before end of line" } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
> +  ;
> +#pragma acc kernels num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
> +  ;
> +#pragma acc kernels vector_length() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
> +  ;
> +
>  #pragma acc parallel num_gangs() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
>    ;
>  #pragma acc parallel num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
> @@ -34,6 +47,14 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs(1 /* { dg-error "expected '\\)' before end of line" } */
> +  ;
> +#pragma acc kernels num_workers(1 /* { dg-error "expected '\\)' before end of line" } */
> +  ;
> +#pragma acc kernels vector_length(1 /* { dg-error "expected '\\)' before end of line" } */
> +  ;
> +
>  #pragma acc parallel num_gangs(1 /* { dg-error "expected '\\)' before end of line" } */
>    ;
>  #pragma acc parallel num_workers(1 /* { dg-error "expected '\\)' before end of line" } */
> @@ -41,6 +62,14 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length(1 /* { dg-error "expected '\\)' before end of line" } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs(i /* { dg-error "expected '\\)' before end of line" } */
> +  ;
> +#pragma acc kernels num_workers(i /* { dg-error "expected '\\)' before end of line" } */
> +  ;
> +#pragma acc kernels vector_length(i /* { dg-error "expected '\\)' before end of line" } */
> +  ;
> +
>  #pragma acc parallel num_gangs(i /* { dg-error "expected '\\)' before end of line" } */
>    ;
>  #pragma acc parallel num_workers(i /* { dg-error "expected '\\)' before end of line" } */
> @@ -48,6 +77,14 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length(i /* { dg-error "expected '\\)' before end of line" } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs(1 i /* { dg-error "expected '\\)' before 'i'" } */
> +  ;
> +#pragma acc kernels num_workers(1 i /* { dg-error "expected '\\)' before 'i'" } */
> +  ;
> +#pragma acc kernels vector_length(1 i /* { dg-error "expected '\\)' before 'i'" } */
> +  ;
> +
>  #pragma acc parallel num_gangs(1 i /* { dg-error "expected '\\)' before 'i'" } */
>    ;
>  #pragma acc parallel num_workers(1 i /* { dg-error "expected '\\)' before 'i'" } */
> @@ -55,6 +92,14 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length(1 i /* { dg-error "expected '\\)' before 'i'" } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs(1 i) /* { dg-error "expected '\\)' before 'i'" } */
> +  ;
> +#pragma acc kernels num_workers(1 i) /* { dg-error "expected '\\)' before 'i'" } */
> +  ;
> +#pragma acc kernels vector_length(1 i) /* { dg-error "expected '\\)' before 'i'" } */
> +  ;
> +
>  #pragma acc parallel num_gangs(1 i) /* { dg-error "expected '\\)' before 'i'" } */
>    ;
>  #pragma acc parallel num_workers(1 i) /* { dg-error "expected '\\)' before 'i'" } */
> @@ -62,6 +107,17 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length(1 i) /* { dg-error "expected '\\)' before 'i'" } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
> +  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
> +  ;
> +#pragma acc kernels num_workers(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
> +  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
> +  ;
> +#pragma acc kernels vector_length(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
> +  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
> +  ;
> +
>  #pragma acc parallel num_gangs(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
>    /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
>    ;
> @@ -72,6 +128,14 @@ void acc_parallel(int i, float f)
>    /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
> +  ;
> +#pragma acc kernels num_workers(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
> +  ;
> +#pragma acc kernels vector_length(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
> +  ;
> +
>  #pragma acc parallel num_gangs(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
>    ;
>  #pragma acc parallel num_workers(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
> @@ -79,11 +143,27 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
>    ;
>  
> -#pragma acc parallel num_gangs(num_gangs) /* { dg-error "'num_gangs' (un|was not )declared" } */
> +
> +#pragma acc kernels num_gangs(num_gangs_k) /* { dg-error "'num_gangs_k' (un|was not )declared" } */
>    ;
> -#pragma acc parallel num_workers(num_workers) /* { dg-error "'num_workers' (un|was not )declared" } */
> +#pragma acc kernels num_workers(num_workers_k) /* { dg-error "'num_workers_k' (un|was not )declared" } */
>    ;
> -#pragma acc parallel vector_length(vector_length) /* { dg-error "'vector_length' (un|was not )declared" } */
> +#pragma acc kernels vector_length(vector_length_k) /* { dg-error "'vector_length_k' (un|was not )declared" } */
> +  ;
> +
> +#pragma acc parallel num_gangs(num_gangs_p) /* { dg-error "'num_gangs_p' (un|was not )declared" } */
> +  ;
> +#pragma acc parallel num_workers(num_workers_p) /* { dg-error "'num_workers_p' (un|was not )declared" } */
> +  ;
> +#pragma acc parallel vector_length(vector_length_p) /* { dg-error "'vector_length_p' (un|was not )declared" } */
> +  ;
> +
> +
> +#pragma acc kernels num_gangs(f) /* { dg-error "'num_gangs' expression must be integral" } */
> +  ;
> +#pragma acc kernels num_workers(f) /* { dg-error "'num_workers' expression must be integral" } */
> +  ;
> +#pragma acc kernels vector_length(f) /* { dg-error "'vector_length' expression must be integral" } */
>    ;
>  
>  #pragma acc parallel num_gangs(f) /* { dg-error "'num_gangs' expression must be integral" } */
> @@ -93,6 +173,14 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length(f) /* { dg-error "'vector_length' expression must be integral" } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs((float) 1) /* { dg-error "'num_gangs' expression must be integral" } */
> +  ;
> +#pragma acc kernels num_workers((float) 1) /* { dg-error "'num_workers' expression must be integral" } */
> +  ;
> +#pragma acc kernels vector_length((float) 1) /* { dg-error "'vector_length' expression must be integral" } */
> +  ;
> +
>  #pragma acc parallel num_gangs((float) 1) /* { dg-error "'num_gangs' expression must be integral" } */
>    ;
>  #pragma acc parallel num_workers((float) 1) /* { dg-error "'num_workers' expression must be integral" } */
> @@ -100,6 +188,14 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length((float) 1) /* { dg-error "'vector_length' expression must be integral" } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs(0) /* { dg-warning "'num_gangs' value must be positive" } */
> +  ;
> +#pragma acc kernels num_workers(0) /* { dg-warning "'num_workers' value must be positive" } */
> +  ;
> +#pragma acc kernels vector_length(0) /* { dg-warning "'vector_length' value must be positive" } */
> +  ;
> +
>  #pragma acc parallel num_gangs(0) /* { dg-warning "'num_gangs' value must be positive" } */
>    ;
>  #pragma acc parallel num_workers(0) /* { dg-warning "'num_workers' value must be positive" } */
> @@ -107,6 +203,14 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length(0) /* { dg-warning "'vector_length' value must be positive" } */
>    ;
>  
> +
> +#pragma acc kernels num_gangs((int) -1.2) /* { dg-warning "'num_gangs' value must be positive" } */
> +  ;
> +#pragma acc kernels num_workers((int) -1.2) /* { dg-warning "'num_workers' value must be positive" } */
> +  ;
> +#pragma acc kernels vector_length((int) -1.2) /* { dg-warning "'vector_length' value must be positive" } */
> +  ;
> +
>  #pragma acc parallel num_gangs((int) -1.2) /* { dg-warning "'num_gangs' value must be positive" } */
>    ;
>  #pragma acc parallel num_workers((int) -1.2) /* { dg-warning "'num_workers' value must be positive" } */
> @@ -114,7 +218,8 @@ void acc_parallel(int i, float f)
>  #pragma acc parallel vector_length((int) -1.2) /* { dg-warning "'vector_length' value must be positive" } */
>    ;
>  
> -#pragma acc parallel \
> +
> +#pragma acc kernels \
>    num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c } } */ \
>    num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c } } */ \
>    vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c } } */ \
> @@ -123,12 +228,31 @@ void acc_parallel(int i, float f)
>    num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c++ } } */
>    ;
>  
> -#pragma acc parallel \
> +#pragma acc parallel							\
> +  num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c } } */ \
> +  num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c } } */ \
> +  vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c } } */ \
> +  num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c++ } } */ \
> +  vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c++ } } */ \
> +  num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c++ } } */
> +  ;
> +
> +
> +#pragma acc kernels \
> +  num_gangs(-1) /* { dg-warning "'num_gangs' value must be positive" } */ \
> +  num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */ \
> +  vector_length(abc_k) /* { dg-error "'abc_k' (un|was not )declared" } */ \
> +  num_workers(0.5) /* { dg-error "'num_workers' expression must be integral" } */ \
> +  vector_length(&f) /* { dg-error "'vector_length' expression must be integral" } */ \
> +  num_gangs( /* { dg-error "expected (primary-|)expression before end of line" "TODO" { xfail c } } */
> +  ;
> +
> +#pragma acc parallel							\
>    num_gangs(-1) /* { dg-warning "'num_gangs' value must be positive" } */ \
>    num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */ \
> -  vector_length(abc) /* { dg-error "'abc' (un|was not )declared" } */ \
> +  vector_length(abc_p) /* { dg-error "'abc_p' (un|was not )declared" } */ \
>    num_workers(0.5) /* { dg-error "'num_workers' expression must be integral" } */ \
> -  vector_length(&acc_parallel) /* { dg-error "'vector_length' expression must be integral" } */ \
> +  vector_length(&f) /* { dg-error "'vector_length' expression must be integral" } */ \
>    num_gangs( /* { dg-error "expected (primary-|)expression before end of line" "TODO" { xfail c } } */
>    ;
>  }
> diff --git gcc/testsuite/c-c++-common/goacc/routine-1.c gcc/testsuite/c-c++-common/goacc/routine-1.c
> index a5e0d69..a756922 100644
> --- gcc/testsuite/c-c++-common/goacc/routine-1.c
> +++ gcc/testsuite/c-c++-common/goacc/routine-1.c
> @@ -21,6 +21,13 @@ void seq (void)
>  
>  int main ()
>  {
> +#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32)
> +  {
> +    gang ();
> +    worker ();
> +    vector ();
> +    seq ();
> +  }
>  
>  #pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
>    {
> diff --git gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
> index 0a006e3..9f11196 100644
> --- gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
> +++ gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
> @@ -1,10 +1,6 @@
> -/* { dg-do compile } */
>  /* { dg-additional-options "-Wuninitialized" } */
>  
> -#include <stdbool.h>
> -
> -int
> -main (void)
> +void acc_parallel()
>  {
>    int i, j, k;
>  
> @@ -17,3 +13,17 @@ main (void)
>    #pragma acc parallel vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
>    ;
>  }
> +
> +void acc_kernels()
> +{
> +  int i, j, k;
> +
> +  #pragma acc kernels num_gangs(i) /* { dg-warning "is used uninitialized in this function" } */
> +  ;
> +
> +  #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized in this function" } */
> +  ;
> +
> +  #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
> +  ;
> +}
> diff --git gcc/testsuite/g++.dg/goacc/template.C gcc/testsuite/g++.dg/goacc/template.C
> index 74f40d8..852f42f 100644
> --- gcc/testsuite/g++.dg/goacc/template.C
> +++ gcc/testsuite/g++.dg/goacc/template.C
> @@ -100,6 +100,10 @@ oacc_kernels_copy (T a)
>    float y = 3;
>    double z = 4;
>  
> +#pragma acc kernels num_gangs (a) num_workers (a) vector_length (a) default (none) copyout (b) copyin (a)
> +  for (int i = 0; i < 1; i++)
> +    b = a;
> +
>  #pragma acc kernels copy (w, x, y, z)
>    {
>      w = accDouble<char>(w);
> diff --git gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
> index 4ec66de..7daca59 100644
> --- gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
> +++ gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
> @@ -6,7 +6,8 @@ program test
>    integer :: q, i, j, k, m, n, o, p, r, s, t, u, v, w
>    logical :: l = .true.
>  
> -  !$acc kernels if(l) async copy(i), copyin(j), copyout(k), create(m) &
> +  !$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
> +  !$acc copy(i), copyin(j), copyout(k), create(m) &
>    !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
>    !$acc deviceptr(u)
>    !$acc end kernels
> @@ -16,6 +17,9 @@ end program test
>  
>  ! { dg-final { scan-tree-dump-times "if" 1 "original" } }
>  ! { dg-final { scan-tree-dump-times "async" 1 "original" } } 
> +! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } 
> +! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } 
> +! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } 
>  
>  ! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } 
>  ! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } 
> diff --git gcc/testsuite/gfortran.dg/goacc/routine-3.f90 gcc/testsuite/gfortran.dg/goacc/routine-3.f90
> index ca9b928..6773f62 100644
> --- gcc/testsuite/gfortran.dg/goacc/routine-3.f90
> +++ gcc/testsuite/gfortran.dg/goacc/routine-3.f90
> @@ -4,6 +4,12 @@ CONTAINS
>      INTEGER  :: i
>      REAL(KIND=8), ALLOCATABLE :: un(:),  ua(:)
>  
> +    !$acc kernels num_gangs(2) num_workers(4) vector_length(32)
> +    DO jj = 1, 100
> +       un(i) = ua(i)
> +    END DO
> +    !$acc end kernels
> +
>      !$acc parallel num_gangs(2) num_workers(4) vector_length(32)
>      DO jj = 1, 100
>         un(i) = ua(i)
> diff --git gcc/testsuite/gfortran.dg/goacc/sie.f95 gcc/testsuite/gfortran.dg/goacc/sie.f95
> index 2d66026..abfe28b 100644
> --- gcc/testsuite/gfortran.dg/goacc/sie.f95
> +++ gcc/testsuite/gfortran.dg/goacc/sie.f95
> @@ -95,6 +95,34 @@ program test
>    !$acc parallel num_gangs("1") ! { dg-error "scalar INTEGER expression" }
>    !$acc end parallel
>  
> +  !$acc kernels num_gangs ! { dg-error "Unclassifiable OpenACC directive" }
> +
> +  !$acc kernels num_gangs(3)
> +  !$acc end kernels
> +
> +  !$acc kernels num_gangs(i)
> +  !$acc end kernels
> +
> +  !$acc kernels num_gangs(i+1)
> +  !$acc end kernels
> +
> +  !$acc kernels num_gangs(-1) ! { dg-warning "must be positive" }
> +  !$acc end kernels
> +
> +  !$acc kernels num_gangs(0) ! { dg-warning "must be positive" }
> +  !$acc end kernels
> +
> +  !$acc kernels num_gangs() ! { dg-error "Invalid character in name" }
> +
> +  !$acc kernels num_gangs(1.5) ! { dg-error "scalar INTEGER expression" }
> +  !$acc end kernels
> +
> +  !$acc kernels num_gangs(.true.) ! { dg-error "scalar INTEGER expression" }
> +  !$acc end kernels
> +
> +  !$acc kernels num_gangs("1") ! { dg-error "scalar INTEGER expression" }
> +  !$acc end kernels
> +
>  
>    !$acc parallel num_workers ! { dg-error "Unclassifiable OpenACC directive" }
>  
> @@ -124,6 +152,34 @@ program test
>    !$acc parallel num_workers("1") ! { dg-error "scalar INTEGER expression" }
>    !$acc end parallel
>  
> +  !$acc kernels num_workers ! { dg-error "Unclassifiable OpenACC directive" }
> +
> +  !$acc kernels num_workers(3)
> +  !$acc end kernels
> +
> +  !$acc kernels num_workers(i)
> +  !$acc end kernels
> +
> +  !$acc kernels num_workers(i+1)
> +  !$acc end kernels
> +
> +  !$acc kernels num_workers(-1) ! { dg-warning "must be positive" }
> +  !$acc end kernels
> +
> +  !$acc kernels num_workers(0) ! { dg-warning "must be positive" }
> +  !$acc end kernels
> +
> +  !$acc kernels num_workers() ! { dg-error "Invalid character in name" }
> +
> +  !$acc kernels num_workers(1.5) ! { dg-error "scalar INTEGER expression" }
> +  !$acc end kernels
> +
> +  !$acc kernels num_workers(.true.) ! { dg-error "scalar INTEGER expression" }
> +  !$acc end kernels
> +
> +  !$acc kernels num_workers("1") ! { dg-error "scalar INTEGER expression" }
> +  !$acc end kernels
> +
>  
>    !$acc parallel vector_length ! { dg-error "Unclassifiable OpenACC directive" }
>  
> @@ -153,6 +209,34 @@ program test
>    !$acc parallel vector_length("1") ! { dg-error "scalar INTEGER expression" }
>    !$acc end parallel
>  
> +  !$acc kernels vector_length ! { dg-error "Unclassifiable OpenACC directive" }
> +
> +  !$acc kernels vector_length(3)
> +  !$acc end kernels
> +
> +  !$acc kernels vector_length(i)
> +  !$acc end kernels
> +
> +  !$acc kernels vector_length(i+1)
> +  !$acc end kernels
> +
> +  !$acc kernels vector_length(-1) ! { dg-warning "must be positive" }
> +  !$acc end kernels
> +
> +  !$acc kernels vector_length(0) ! { dg-warning "must be positive" }
> +  !$acc end kernels
> +
> +  !$acc kernels vector_length() ! { dg-error "Invalid character in name" }
> +
> +  !$acc kernels vector_length(1.5) ! { dg-error "scalar INTEGER expression" }
> +  !$acc end kernels
> +
> +  !$acc kernels vector_length(.true.) ! { dg-error "scalar INTEGER expression" }
> +  !$acc end kernels
> +
> +  !$acc kernels vector_length("1") ! { dg-error "scalar INTEGER expression" }
> +  !$acc end kernels
> +
>  
>    !$acc loop gang
>    do i = 1,10
> @@ -249,4 +333,4 @@ program test
>    do i = 1,10
>    enddo
>  
> -end program test
> \ No newline at end of file
> +end program test
> diff --git gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
> index b87d26f..5dea42b 100644
> --- gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
> +++ gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
> @@ -1,7 +1,6 @@
> -! { dg-do compile }
>  ! { dg-additional-options "-Wuninitialized" }
>  
> -program test
> +subroutine acc_parallel
>    implicit none
>    integer :: i, j, k
>  
> @@ -13,5 +12,18 @@ program test
>  
>    !$acc parallel vector_length(k) ! { dg-warning "is used uninitialized in this function" }
>    !$acc end parallel
> +end subroutine acc_parallel
>  
> -end program test
> +subroutine acc_kernels
> +  implicit none
> +  integer :: i, j, k
> +
> +  !$acc kernels num_gangs(i) ! { dg-warning "is used uninitialized in this function" }
> +  !$acc end kernels
> +
> +  !$acc kernels num_workers(j) ! { dg-warning "is used uninitialized in this function" }
> +  !$acc end kernels
> +
> +  !$acc kernels vector_length(k) ! { dg-warning "is used uninitialized in this function" }
> +  !$acc end kernels
> +end subroutine acc_kernels
> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
> index c7592d6..b840888 100644
> --- libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
> +++ libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
> @@ -14,27 +14,40 @@ main (void)
>    b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
>    c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
>  
> +  /* Parallelism dimensions: compiler/runtime decides.  */
>  #pragma acc kernels copyout (a[0:N])
>    {
>      for (COUNTERTYPE i = 0; i < N; i++)
>        a[i] = i * 2;
>    }
>  
> -#pragma acc kernels copyout (b[0:N])
> +  /* Parallelism dimensions: variable.  */
> +#pragma acc kernels copyout (b[0:N]) \
> +  num_gangs (3 + a[3]) num_workers (5 + a[5]) vector_length (7 + a[7])
> +  /* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */
>    {
>      for (COUNTERTYPE i = 0; i < N; i++)
>        b[i] = i * 4;
>    }
>  
> -#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
> +  /* Parallelism dimensions: literal.  */
> +#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) \
> +  num_gangs (3) num_workers (5) vector_length (7)
> +  /* { dg-prune-output "using vector_length \\(32\\), ignoring 7" } */
>    {
>      for (COUNTERTYPE ii = 0; ii < N; ii++)
>        c[ii] = a[ii] + b[ii];
>    }
>  
>    for (COUNTERTYPE i = 0; i < N; i++)
> -    if (c[i] != a[i] + b[i])
> -      abort ();
> +    {
> +      if (a[i] != i * 2)
> +	abort ();
> +      if (b[i] != i * 4)
> +	abort ();
> +      if (c[i] != a[i] + b[i])
> +	abort ();
> +    }
>  
>    free (a);
>    free (b);
> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
> index d8af546..8308f7c 100644
> --- libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
> +++ libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
> @@ -520,5 +520,40 @@ int main ()
>    }
>  
>  
> +  /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1
> +     kernels even when there are explicit num_gangs, num_workers, or
> +     vector_length clauses.  */
> +  {
> +    int gangs = 5;
> +#define WORKERS 5
> +#define VECTORS 13
> +    int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
> +    gangs_min = workers_min = vectors_min = INT_MAX;
> +    gangs_max = workers_max = vectors_max = INT_MIN;
> +#pragma acc kernels \
> +  num_gangs (gangs) \
> +  num_workers (WORKERS) \
> +  vector_length (VECTORS)
> +    {
> +      /* This is to make the OpenACC kernels construct unparallelizable.  */
> +      asm volatile ("" : : : "memory");
> +
> +#pragma acc loop reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
> +      for (int i = 100; i > -100; --i)
> +	{
> +	  gangs_min = gangs_max = acc_gang ();
> +	  workers_min = workers_max = acc_worker ();
> +	  vectors_min = vectors_max = acc_vector ();
> +	}
> +    }
> +    if (gangs_min != 0 || gangs_max != 1 - 1
> +	|| workers_min != 0 || workers_max != 1 - 1
> +	|| vectors_min != 0 || vectors_max != 1 - 1)
> +      __builtin_abort ();
> +#undef VECTORS
> +#undef WORKERS
> +  }
> +
> +
>    return 0;
>  }
> diff --git libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95 libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
> index 163e8d5..b88ca67 100644
> --- libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
> +++ libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
> @@ -6,25 +6,34 @@ program main
>    integer, dimension (0:n-1) :: a, b, c
>    integer                    :: i, ii
>  
> +  ! Parallelism dimensions: compiler/runtime decides.
>    !$acc kernels copyout (a(0:n-1))
>    do i = 0, n - 1
>       a(i) = i * 2
>    end do
>    !$acc end kernels
>  
> -  !$acc kernels copyout (b(0:n-1))
> +  ! Parallelism dimensions: variable.
> +  !$acc kernels copyout (b(0:n-1)) &
> +  !$acc num_gangs (3 + a(3)) num_workers (5 + a(5)) vector_length (7 + a(7))
> +  ! { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" }
>    do i = 0, n -1
>       b(i) = i * 4
>    end do
>    !$acc end kernels
>  
> -  !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
> +  ! Parallelism dimensions: literal.
> +  !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) &
> +  !$acc num_gangs (3) num_workers (5) vector_length (7)
> +  ! { dg-prune-output "using vector_length \\(32\\), ignoring 7" }
>    do ii = 0, n - 1
>       c(ii) = a(ii) + b(ii)
>    end do
>    !$acc end kernels
>  
>    do i = 0, n - 1
> +     if (a(i) .ne. i * 2) call abort
> +     if (b(i) .ne. i * 4) call abort
>       if (c(i) .ne. a(i) + b(i)) call abort
>    end do


Grüße
 Thomas

^ permalink raw reply	[flat|nested] 5+ messages in thread

* Re: OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses
  2017-05-11 12:28 OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses Thomas Schwinge
  2017-05-14 10:40 ` Thomas Schwinge
  2017-05-19 11:09 ` Thomas Schwinge
@ 2017-05-23  9:54 ` Jakub Jelinek
  2017-05-23 15:58   ` Thomas Schwinge
  2 siblings, 1 reply; 5+ messages in thread
From: Jakub Jelinek @ 2017-05-23  9:54 UTC (permalink / raw)
  To: Thomas Schwinge; +Cc: gcc-patches

On Thu, May 11, 2017 at 02:26:51PM +0200, Thomas Schwinge wrote:
>     OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses
>     
>             gcc/c/
>             * c-parser.c (OACC_KERNELS_CLAUSE_MASK): Add
>             "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
>             "VECTOR_LENGTH".
>             gcc/cp/
>             * parser.c (OACC_KERNELS_CLAUSE_MASK): Add
>             "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
>             "VECTOR_LENGTH".
>             gcc/fortran/
>             * openmp.c (OACC_KERNELS_CLAUSES): Add "OMP_CLAUSE_NUM_GANGS",
>             "OMP_CLAUSE_NUM_WORKERS", "OMP_CLAUSE_VECTOR_LENGTH".
>             gcc/
>             * omp-offload.c (execute_oacc_device_lower): Remove the
>             parallelism dimensions function attributes for unparallelized
>             OpenACC kernels constructs.
>             gcc/testsuite/
>             * c-c++-common/goacc/parallel-dims-1.c: Update.
>             * c-c++-common/goacc/parallel-dims-2.c: Likewise.
>             * c-c++-common/goacc/routine-1.c: Likewise.
>             * c-c++-common/goacc/uninit-dim-clause.c: Likewise.
>             * g++.dg/goacc/template.C: Likewise.
>             * gfortran.dg/goacc/kernels-tree.f95: Likewise.
>             * gfortran.dg/goacc/routine-3.f90: Likewise.
>             * gfortran.dg/goacc/sie.f95: Likewise.
>             * gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
>             libgomp/
>             * testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: Update.
>             * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
>             * testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise.

Ok, thanks.

	Jakub

^ permalink raw reply	[flat|nested] 5+ messages in thread

* Re: OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses
  2017-05-23  9:54 ` Jakub Jelinek
@ 2017-05-23 15:58   ` Thomas Schwinge
  0 siblings, 0 replies; 5+ messages in thread
From: Thomas Schwinge @ 2017-05-23 15:58 UTC (permalink / raw)
  To: Jakub Jelinek, gcc-patches

Hi!

On Tue, 23 May 2017 11:48:09 +0200, Jakub Jelinek <jakub@redhat.com> wrote:
> On Thu, May 11, 2017 at 02:26:51PM +0200, Thomas Schwinge wrote:
> >     OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses

> Ok, thanks.

Thanks.  As posted, committed to trunk in r248370:

commit 9d5c2cca06bf15e6eff22c3a4f6e1cf0072645e5
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Tue May 23 15:47:32 2017 +0000

    OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses
    
            gcc/c/
            * c-parser.c (OACC_KERNELS_CLAUSE_MASK): Add
            "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
            "VECTOR_LENGTH".
            gcc/cp/
            * parser.c (OACC_KERNELS_CLAUSE_MASK): Add
            "PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
            "VECTOR_LENGTH".
            gcc/fortran/
            * openmp.c (OACC_KERNELS_CLAUSES): Add "OMP_CLAUSE_NUM_GANGS",
            "OMP_CLAUSE_NUM_WORKERS", "OMP_CLAUSE_VECTOR_LENGTH".
            gcc/
            * omp-offload.c (execute_oacc_device_lower): Remove the
            parallelism dimensions function attributes for unparallelized
            OpenACC kernels constructs.
            gcc/testsuite/
            * c-c++-common/goacc/parallel-dims-1.c: Update.
            * c-c++-common/goacc/parallel-dims-2.c: Likewise.
            * c-c++-common/goacc/routine-1.c: Likewise.
            * c-c++-common/goacc/uninit-dim-clause.c: Likewise.
            * g++.dg/goacc/template.C: Likewise.
            * gfortran.dg/goacc/kernels-tree.f95: Likewise.
            * gfortran.dg/goacc/routine-3.f90: Likewise.
            * gfortran.dg/goacc/sie.f95: Likewise.
            * gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
            libgomp/
            * testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: Update.
            * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
            * testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@248370 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog                                      |   6 +
 gcc/c/ChangeLog                                    |   6 +
 gcc/c/c-parser.c                                   |   3 +
 gcc/cp/ChangeLog                                   |   6 +
 gcc/cp/parser.c                                    |   3 +
 gcc/fortran/ChangeLog                              |   5 +
 gcc/fortran/openmp.c                               |   3 +-
 gcc/omp-offload.c                                  |   9 ++
 gcc/testsuite/ChangeLog                            |  12 ++
 gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c |   3 +
 gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c | 152 +++++++++++++++++++--
 gcc/testsuite/c-c++-common/goacc/routine-1.c       |   7 +
 .../c-c++-common/goacc/uninit-dim-clause.c         |  20 ++-
 gcc/testsuite/g++.dg/goacc/template.C              |   4 +
 gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95   |   6 +-
 gcc/testsuite/gfortran.dg/goacc/routine-3.f90      |   6 +
 gcc/testsuite/gfortran.dg/goacc/sie.f95            |  86 +++++++++++-
 .../gfortran.dg/goacc/uninit-dim-clause.f95        |  18 ++-
 libgomp/ChangeLog                                  |   4 +
 .../libgomp.oacc-c-c++-common/kernels-loop-2.c     |  21 ++-
 .../libgomp.oacc-c-c++-common/parallel-dims.c      |  35 +++++
 .../libgomp.oacc-fortran/kernels-loop-2.f95        |  13 +-
 22 files changed, 397 insertions(+), 31 deletions(-)

diff --git gcc/ChangeLog gcc/ChangeLog
index d2e846e..b38a31d 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,9 @@
+2017-05-23  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* omp-offload.c (execute_oacc_device_lower): Remove the
+	parallelism dimensions function attributes for unparallelized
+	OpenACC kernels constructs.
+
 2017-05-23  Martin Liska  <mliska@suse.cz>
 
 	* cgraph.c (cgraph_node::get_create): Use symtab_node::dump_{asm_,}name
diff --git gcc/c/ChangeLog gcc/c/ChangeLog
index d768d93..cb04d4a 100644
--- gcc/c/ChangeLog
+++ gcc/c/ChangeLog
@@ -1,3 +1,9 @@
+2017-05-23  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-parser.c (OACC_KERNELS_CLAUSE_MASK): Add
+	"PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
+	"VECTOR_LENGTH".
+
 2017-05-23  Marek Polacek  <polacek@redhat.com>
 
 	* c-parser.c (c_parser_compound_statement_nostart): Remove redundant
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index f3bcbee..03c711b 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -13984,11 +13984,14 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
 #define OACC_PARALLEL_CLAUSE_MASK					\
diff --git gcc/cp/ChangeLog gcc/cp/ChangeLog
index aaf39c2..bfe718f 100644
--- gcc/cp/ChangeLog
+++ gcc/cp/ChangeLog
@@ -1,3 +1,9 @@
+2017-05-23  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* parser.c (OACC_KERNELS_CLAUSE_MASK): Add
+	"PRAGMA_OACC_CLAUSE_NUM_GANGS", "PRAGMA_OACC_CLAUSE_NUM_WORKERS",
+	"VECTOR_LENGTH".
+
 2017-05-23  Nathan Sidwell  <nathan@acm.org>
 
 	* cp-tree.h (OVL_P): New.
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 23d979c..b39e624 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -36432,11 +36432,14 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)			\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS)		\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT)		\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE)	\
+	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH)	\
 	| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
 
 #define OACC_PARALLEL_CLAUSE_MASK					\
diff --git gcc/fortran/ChangeLog gcc/fortran/ChangeLog
index 6977bd1..b3179e0 100644
--- gcc/fortran/ChangeLog
+++ gcc/fortran/ChangeLog
@@ -1,3 +1,8 @@
+2017-05-23  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* openmp.c (OACC_KERNELS_CLAUSES): Add "OMP_CLAUSE_NUM_GANGS",
+	"OMP_CLAUSE_NUM_WORKERS", "OMP_CLAUSE_VECTOR_LENGTH".
+
 2017-05-22  Janus Weil  <janus@gcc.gnu.org>
 
 	PR fortran/80766
diff --git gcc/fortran/openmp.c gcc/fortran/openmp.c
index 80146e2..5a2b774 100644
--- gcc/fortran/openmp.c
+++ gcc/fortran/openmp.c
@@ -1932,7 +1932,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
    | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE \
    | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
 #define OACC_KERNELS_CLAUSES \
-  (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_DEVICEPTR	      \
+  (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS	      \
+   | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
    | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT		      \
    | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY      \
    | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT	      \
diff --git gcc/omp-offload.c gcc/omp-offload.c
index f02b4f8..54a4e90 100644
--- gcc/omp-offload.c
+++ gcc/omp-offload.c
@@ -1451,6 +1451,15 @@ execute_oacc_device_lower ()
     = (lookup_attribute ("oacc kernels parallelized",
 			 DECL_ATTRIBUTES (current_function_decl)) != NULL);
 
+  /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1
+     kernels, so remove the parallelism dimensions function attributes
+     potentially set earlier on.  */
+  if (is_oacc_kernels && !is_oacc_kernels_parallelized)
+    {
+      oacc_set_fn_attrib (current_function_decl, NULL, NULL);
+      attrs = oacc_get_fn_attrib (current_function_decl);
+    }
+
   /* Discover, partition and process the loops.  */
   oacc_loop *loops = oacc_loop_discovery ();
   int fn_level = oacc_fn_attrib_level (attrs);
diff --git gcc/testsuite/ChangeLog gcc/testsuite/ChangeLog
index 7b42abc..cef3ba6 100644
--- gcc/testsuite/ChangeLog
+++ gcc/testsuite/ChangeLog
@@ -1,3 +1,15 @@
+2017-05-23  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/parallel-dims-1.c: Update.
+	* c-c++-common/goacc/parallel-dims-2.c: Likewise.
+	* c-c++-common/goacc/routine-1.c: Likewise.
+	* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
+	* g++.dg/goacc/template.C: Likewise.
+	* gfortran.dg/goacc/kernels-tree.f95: Likewise.
+	* gfortran.dg/goacc/routine-3.f90: Likewise.
+	* gfortran.dg/goacc/sie.f95: Likewise.
+	* gfortran.dg/goacc/uninit-dim-clause.f95: Likewise.
+
 2017-05-23  Nathan Sidwell  <nathan@acm.org>
 
 	* g++.dg/lookup/using13.C: Adjust expected error.
diff --git gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
index a85d3d3..57f682f 100644
--- gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
+++ gcc/testsuite/c-c++-common/goacc/parallel-dims-1.c
@@ -3,6 +3,9 @@
 
 void f(int i)
 {
+#pragma acc kernels num_gangs(i) num_workers(i) vector_length(i)
+  ;
+
 #pragma acc parallel num_gangs(i) num_workers(i) vector_length(i)
   ;
 }
diff --git gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
index 30a3d17..acfbe7f 100644
--- gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
+++ gcc/testsuite/c-c++-common/goacc/parallel-dims-2.c
@@ -1,18 +1,15 @@
 /* Invalid use of OpenACC parallelism dimensions clauses: num_gangs,
    num_workers, vector_length.  */
 
-void acc_kernels(int i)
+void f(int i, float f)
 {
-#pragma acc kernels num_gangs(i) /* { dg-error "'num_gangs' is not valid for '#pragma acc kernels'" } */
+#pragma acc kernels num_gangs /* { dg-error "expected '\\(' before end of line" } */
   ;
-#pragma acc kernels num_workers(i) /* { dg-error "'num_workers' is not valid for '#pragma acc kernels'" } */
+#pragma acc kernels num_workers /* { dg-error "expected '\\(' before end of line" } */
   ;
-#pragma acc kernels vector_length(i) /* { dg-error "'vector_length' is not valid for '#pragma acc kernels'" } */
+#pragma acc kernels vector_length /* { dg-error "expected '\\(' before end of line" } */
   ;
-}
 
-void acc_parallel(int i, float f)
-{
 #pragma acc parallel num_gangs /* { dg-error "expected '\\(' before end of line" } */
   ;
 #pragma acc parallel num_workers /* { dg-error "expected '\\(' before end of line" } */
@@ -20,6 +17,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length /* { dg-error "expected '\\(' before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs( /* { dg-error "expected (primary-|)expression before end of line" } */
+  ;
+#pragma acc kernels num_workers( /* { dg-error "expected (primary-|)expression before end of line" } */
+  ;
+#pragma acc kernels vector_length( /* { dg-error "expected (primary-|)expression before end of line" } */
+  ;
+
 #pragma acc parallel num_gangs( /* { dg-error "expected (primary-|)expression before end of line" } */
   ;
 #pragma acc parallel num_workers( /* { dg-error "expected (primary-|)expression before end of line" } */
@@ -27,6 +32,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length( /* { dg-error "expected (primary-|)expression before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
+  ;
+#pragma acc kernels num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
+  ;
+#pragma acc kernels vector_length() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
+  ;
+
 #pragma acc parallel num_gangs() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
   ;
 #pragma acc parallel num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
@@ -34,6 +47,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length() /* { dg-error "expected (primary-|)expression before '\\)' token" } */
   ;
 
+
+#pragma acc kernels num_gangs(1 /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels num_workers(1 /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels vector_length(1 /* { dg-error "expected '\\)' before end of line" } */
+  ;
+
 #pragma acc parallel num_gangs(1 /* { dg-error "expected '\\)' before end of line" } */
   ;
 #pragma acc parallel num_workers(1 /* { dg-error "expected '\\)' before end of line" } */
@@ -41,6 +62,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1 /* { dg-error "expected '\\)' before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs(i /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels num_workers(i /* { dg-error "expected '\\)' before end of line" } */
+  ;
+#pragma acc kernels vector_length(i /* { dg-error "expected '\\)' before end of line" } */
+  ;
+
 #pragma acc parallel num_gangs(i /* { dg-error "expected '\\)' before end of line" } */
   ;
 #pragma acc parallel num_workers(i /* { dg-error "expected '\\)' before end of line" } */
@@ -48,6 +77,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(i /* { dg-error "expected '\\)' before end of line" } */
   ;
 
+
+#pragma acc kernels num_gangs(1 i /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels num_workers(1 i /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels vector_length(1 i /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+
 #pragma acc parallel num_gangs(1 i /* { dg-error "expected '\\)' before 'i'" } */
   ;
 #pragma acc parallel num_workers(1 i /* { dg-error "expected '\\)' before 'i'" } */
@@ -55,6 +92,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1 i /* { dg-error "expected '\\)' before 'i'" } */
   ;
 
+
+#pragma acc kernels num_gangs(1 i) /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels num_workers(1 i) /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+#pragma acc kernels vector_length(1 i) /* { dg-error "expected '\\)' before 'i'" } */
+  ;
+
 #pragma acc parallel num_gangs(1 i) /* { dg-error "expected '\\)' before 'i'" } */
   ;
 #pragma acc parallel num_workers(1 i) /* { dg-error "expected '\\)' before 'i'" } */
@@ -62,6 +107,17 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1 i) /* { dg-error "expected '\\)' before 'i'" } */
   ;
 
+
+#pragma acc kernels num_gangs(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
+  ;
+#pragma acc kernels num_workers(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
+  ;
+#pragma acc kernels vector_length(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
+  ;
+
 #pragma acc parallel num_gangs(1, i /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
   /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
   ;
@@ -72,6 +128,14 @@ void acc_parallel(int i, float f)
   /* { dg-bogus "expected '\\)' before end of line" "TODO" { xfail c } .-1 } */
   ;
 
+
+#pragma acc kernels num_gangs(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  ;
+#pragma acc kernels num_workers(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  ;
+#pragma acc kernels vector_length(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
+  ;
+
 #pragma acc parallel num_gangs(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
   ;
 #pragma acc parallel num_workers(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
@@ -79,11 +143,27 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(1, i) /* { dg-error "expected '\\)' before ',' token" "TODO" { xfail c } } */
   ;
 
-#pragma acc parallel num_gangs(num_gangs) /* { dg-error "'num_gangs' (un|was not )declared" } */
+
+#pragma acc kernels num_gangs(num_gangs_k) /* { dg-error "'num_gangs_k' (un|was not )declared" } */
   ;
-#pragma acc parallel num_workers(num_workers) /* { dg-error "'num_workers' (un|was not )declared" } */
+#pragma acc kernels num_workers(num_workers_k) /* { dg-error "'num_workers_k' (un|was not )declared" } */
   ;
-#pragma acc parallel vector_length(vector_length) /* { dg-error "'vector_length' (un|was not )declared" } */
+#pragma acc kernels vector_length(vector_length_k) /* { dg-error "'vector_length_k' (un|was not )declared" } */
+  ;
+
+#pragma acc parallel num_gangs(num_gangs_p) /* { dg-error "'num_gangs_p' (un|was not )declared" } */
+  ;
+#pragma acc parallel num_workers(num_workers_p) /* { dg-error "'num_workers_p' (un|was not )declared" } */
+  ;
+#pragma acc parallel vector_length(vector_length_p) /* { dg-error "'vector_length_p' (un|was not )declared" } */
+  ;
+
+
+#pragma acc kernels num_gangs(f) /* { dg-error "'num_gangs' expression must be integral" } */
+  ;
+#pragma acc kernels num_workers(f) /* { dg-error "'num_workers' expression must be integral" } */
+  ;
+#pragma acc kernels vector_length(f) /* { dg-error "'vector_length' expression must be integral" } */
   ;
 
 #pragma acc parallel num_gangs(f) /* { dg-error "'num_gangs' expression must be integral" } */
@@ -93,6 +173,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(f) /* { dg-error "'vector_length' expression must be integral" } */
   ;
 
+
+#pragma acc kernels num_gangs((float) 1) /* { dg-error "'num_gangs' expression must be integral" } */
+  ;
+#pragma acc kernels num_workers((float) 1) /* { dg-error "'num_workers' expression must be integral" } */
+  ;
+#pragma acc kernels vector_length((float) 1) /* { dg-error "'vector_length' expression must be integral" } */
+  ;
+
 #pragma acc parallel num_gangs((float) 1) /* { dg-error "'num_gangs' expression must be integral" } */
   ;
 #pragma acc parallel num_workers((float) 1) /* { dg-error "'num_workers' expression must be integral" } */
@@ -100,6 +188,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length((float) 1) /* { dg-error "'vector_length' expression must be integral" } */
   ;
 
+
+#pragma acc kernels num_gangs(0) /* { dg-warning "'num_gangs' value must be positive" } */
+  ;
+#pragma acc kernels num_workers(0) /* { dg-warning "'num_workers' value must be positive" } */
+  ;
+#pragma acc kernels vector_length(0) /* { dg-warning "'vector_length' value must be positive" } */
+  ;
+
 #pragma acc parallel num_gangs(0) /* { dg-warning "'num_gangs' value must be positive" } */
   ;
 #pragma acc parallel num_workers(0) /* { dg-warning "'num_workers' value must be positive" } */
@@ -107,6 +203,14 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length(0) /* { dg-warning "'vector_length' value must be positive" } */
   ;
 
+
+#pragma acc kernels num_gangs((int) -1.2) /* { dg-warning "'num_gangs' value must be positive" } */
+  ;
+#pragma acc kernels num_workers((int) -1.2) /* { dg-warning "'num_workers' value must be positive" } */
+  ;
+#pragma acc kernels vector_length((int) -1.2) /* { dg-warning "'vector_length' value must be positive" } */
+  ;
+
 #pragma acc parallel num_gangs((int) -1.2) /* { dg-warning "'num_gangs' value must be positive" } */
   ;
 #pragma acc parallel num_workers((int) -1.2) /* { dg-warning "'num_workers' value must be positive" } */
@@ -114,7 +218,8 @@ void acc_parallel(int i, float f)
 #pragma acc parallel vector_length((int) -1.2) /* { dg-warning "'vector_length' value must be positive" } */
   ;
 
-#pragma acc parallel \
+
+#pragma acc kernels \
   num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c } } */ \
   num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c } } */ \
   vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c } } */ \
@@ -123,12 +228,31 @@ void acc_parallel(int i, float f)
   num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c++ } } */
   ;
 
-#pragma acc parallel \
+#pragma acc parallel							\
+  num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c } } */ \
+  num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c } } */ \
+  vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c } } */ \
+  num_workers(1) /* { dg-error "too many 'num_workers' clauses" "" { target c++ } } */ \
+  vector_length(1) /* { dg-error "too many 'vector_length' clauses" "" { target c++ } } */ \
+  num_gangs(1) /* { dg-error "too many 'num_gangs' clauses" "" { target c++ } } */
+  ;
+
+
+#pragma acc kernels \
+  num_gangs(-1) /* { dg-warning "'num_gangs' value must be positive" } */ \
+  num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */ \
+  vector_length(abc_k) /* { dg-error "'abc_k' (un|was not )declared" } */ \
+  num_workers(0.5) /* { dg-error "'num_workers' expression must be integral" } */ \
+  vector_length(&f) /* { dg-error "'vector_length' expression must be integral" } */ \
+  num_gangs( /* { dg-error "expected (primary-|)expression before end of line" "TODO" { xfail c } } */
+  ;
+
+#pragma acc parallel							\
   num_gangs(-1) /* { dg-warning "'num_gangs' value must be positive" } */ \
   num_workers() /* { dg-error "expected (primary-|)expression before '\\)' token" } */ \
-  vector_length(abc) /* { dg-error "'abc' (un|was not )declared" } */ \
+  vector_length(abc_p) /* { dg-error "'abc_p' (un|was not )declared" } */ \
   num_workers(0.5) /* { dg-error "'num_workers' expression must be integral" } */ \
-  vector_length(&acc_parallel) /* { dg-error "'vector_length' expression must be integral" } */ \
+  vector_length(&f) /* { dg-error "'vector_length' expression must be integral" } */ \
   num_gangs( /* { dg-error "expected (primary-|)expression before end of line" "TODO" { xfail c } } */
   ;
 }
diff --git gcc/testsuite/c-c++-common/goacc/routine-1.c gcc/testsuite/c-c++-common/goacc/routine-1.c
index a5e0d69..a756922 100644
--- gcc/testsuite/c-c++-common/goacc/routine-1.c
+++ gcc/testsuite/c-c++-common/goacc/routine-1.c
@@ -21,6 +21,13 @@ void seq (void)
 
 int main ()
 {
+#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32)
+  {
+    gang ();
+    worker ();
+    vector ();
+    seq ();
+  }
 
 #pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
   {
diff --git gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
index 0a006e3..9f11196 100644
--- gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
+++ gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
@@ -1,10 +1,6 @@
-/* { dg-do compile } */
 /* { dg-additional-options "-Wuninitialized" } */
 
-#include <stdbool.h>
-
-int
-main (void)
+void acc_parallel()
 {
   int i, j, k;
 
@@ -17,3 +13,17 @@ main (void)
   #pragma acc parallel vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
   ;
 }
+
+void acc_kernels()
+{
+  int i, j, k;
+
+  #pragma acc kernels num_gangs(i) /* { dg-warning "is used uninitialized in this function" } */
+  ;
+
+  #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized in this function" } */
+  ;
+
+  #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
+  ;
+}
diff --git gcc/testsuite/g++.dg/goacc/template.C gcc/testsuite/g++.dg/goacc/template.C
index 74f40d8..852f42f 100644
--- gcc/testsuite/g++.dg/goacc/template.C
+++ gcc/testsuite/g++.dg/goacc/template.C
@@ -100,6 +100,10 @@ oacc_kernels_copy (T a)
   float y = 3;
   double z = 4;
 
+#pragma acc kernels num_gangs (a) num_workers (a) vector_length (a) default (none) copyout (b) copyin (a)
+  for (int i = 0; i < 1; i++)
+    b = a;
+
 #pragma acc kernels copy (w, x, y, z)
   {
     w = accDouble<char>(w);
diff --git gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
index 4ec66de..7daca59 100644
--- gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -6,7 +6,8 @@ program test
   integer :: q, i, j, k, m, n, o, p, r, s, t, u, v, w
   logical :: l = .true.
 
-  !$acc kernels if(l) async copy(i), copyin(j), copyout(k), create(m) &
+  !$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
+  !$acc copy(i), copyin(j), copyout(k), create(m) &
   !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
   !$acc deviceptr(u)
   !$acc end kernels
@@ -16,6 +17,9 @@ end program test
 
 ! { dg-final { scan-tree-dump-times "if" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "async" 1 "original" } } 
+! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } 
+! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } 
+! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } 
 
 ! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } } 
 ! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } } 
diff --git gcc/testsuite/gfortran.dg/goacc/routine-3.f90 gcc/testsuite/gfortran.dg/goacc/routine-3.f90
index ca9b928..6773f62 100644
--- gcc/testsuite/gfortran.dg/goacc/routine-3.f90
+++ gcc/testsuite/gfortran.dg/goacc/routine-3.f90
@@ -4,6 +4,12 @@ CONTAINS
     INTEGER  :: i
     REAL(KIND=8), ALLOCATABLE :: un(:),  ua(:)
 
+    !$acc kernels num_gangs(2) num_workers(4) vector_length(32)
+    DO jj = 1, 100
+       un(i) = ua(i)
+    END DO
+    !$acc end kernels
+
     !$acc parallel num_gangs(2) num_workers(4) vector_length(32)
     DO jj = 1, 100
        un(i) = ua(i)
diff --git gcc/testsuite/gfortran.dg/goacc/sie.f95 gcc/testsuite/gfortran.dg/goacc/sie.f95
index 2d66026..abfe28b 100644
--- gcc/testsuite/gfortran.dg/goacc/sie.f95
+++ gcc/testsuite/gfortran.dg/goacc/sie.f95
@@ -95,6 +95,34 @@ program test
   !$acc parallel num_gangs("1") ! { dg-error "scalar INTEGER expression" }
   !$acc end parallel
 
+  !$acc kernels num_gangs ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc kernels num_gangs(3)
+  !$acc end kernels
+
+  !$acc kernels num_gangs(i)
+  !$acc end kernels
+
+  !$acc kernels num_gangs(i+1)
+  !$acc end kernels
+
+  !$acc kernels num_gangs(-1) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs(0) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs() ! { dg-error "Invalid character in name" }
+
+  !$acc kernels num_gangs(1.5) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs(.true.) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_gangs("1") ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
 
   !$acc parallel num_workers ! { dg-error "Unclassifiable OpenACC directive" }
 
@@ -124,6 +152,34 @@ program test
   !$acc parallel num_workers("1") ! { dg-error "scalar INTEGER expression" }
   !$acc end parallel
 
+  !$acc kernels num_workers ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc kernels num_workers(3)
+  !$acc end kernels
+
+  !$acc kernels num_workers(i)
+  !$acc end kernels
+
+  !$acc kernels num_workers(i+1)
+  !$acc end kernels
+
+  !$acc kernels num_workers(-1) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_workers(0) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels num_workers() ! { dg-error "Invalid character in name" }
+
+  !$acc kernels num_workers(1.5) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_workers(.true.) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels num_workers("1") ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
 
   !$acc parallel vector_length ! { dg-error "Unclassifiable OpenACC directive" }
 
@@ -153,6 +209,34 @@ program test
   !$acc parallel vector_length("1") ! { dg-error "scalar INTEGER expression" }
   !$acc end parallel
 
+  !$acc kernels vector_length ! { dg-error "Unclassifiable OpenACC directive" }
+
+  !$acc kernels vector_length(3)
+  !$acc end kernels
+
+  !$acc kernels vector_length(i)
+  !$acc end kernels
+
+  !$acc kernels vector_length(i+1)
+  !$acc end kernels
+
+  !$acc kernels vector_length(-1) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels vector_length(0) ! { dg-warning "must be positive" }
+  !$acc end kernels
+
+  !$acc kernels vector_length() ! { dg-error "Invalid character in name" }
+
+  !$acc kernels vector_length(1.5) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels vector_length(.true.) ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
+  !$acc kernels vector_length("1") ! { dg-error "scalar INTEGER expression" }
+  !$acc end kernels
+
 
   !$acc loop gang
   do i = 1,10
@@ -249,4 +333,4 @@ program test
   do i = 1,10
   enddo
 
-end program test
\ No newline at end of file
+end program test
diff --git gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95 gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
index b87d26f..5dea42b 100644
--- gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
+++ gcc/testsuite/gfortran.dg/goacc/uninit-dim-clause.f95
@@ -1,7 +1,6 @@
-! { dg-do compile }
 ! { dg-additional-options "-Wuninitialized" }
 
-program test
+subroutine acc_parallel
   implicit none
   integer :: i, j, k
 
@@ -13,5 +12,18 @@ program test
 
   !$acc parallel vector_length(k) ! { dg-warning "is used uninitialized in this function" }
   !$acc end parallel
+end subroutine acc_parallel
 
-end program test
+subroutine acc_kernels
+  implicit none
+  integer :: i, j, k
+
+  !$acc kernels num_gangs(i) ! { dg-warning "is used uninitialized in this function" }
+  !$acc end kernels
+
+  !$acc kernels num_workers(j) ! { dg-warning "is used uninitialized in this function" }
+  !$acc end kernels
+
+  !$acc kernels vector_length(k) ! { dg-warning "is used uninitialized in this function" }
+  !$acc end kernels
+end subroutine acc_kernels
diff --git libgomp/ChangeLog libgomp/ChangeLog
index 8fd5f07..32f8bf1 100644
--- libgomp/ChangeLog
+++ libgomp/ChangeLog
@@ -1,5 +1,9 @@
 2017-05-23  Thomas Schwinge  <thomas@codesourcery.com>
 
+	* testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: Update.
+	* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
+	* testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise.
+
 	* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Rewrite.
 	* testsuite/lib/libgomp.exp
 	(check_effective_target_openacc_nvidia_accel_configured): New
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
index c7592d6..b840888 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c
@@ -14,27 +14,40 @@ main (void)
   b = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
   c = (unsigned int *__restrict)malloc (N * sizeof (unsigned int));
 
+  /* Parallelism dimensions: compiler/runtime decides.  */
 #pragma acc kernels copyout (a[0:N])
   {
     for (COUNTERTYPE i = 0; i < N; i++)
       a[i] = i * 2;
   }
 
-#pragma acc kernels copyout (b[0:N])
+  /* Parallelism dimensions: variable.  */
+#pragma acc kernels copyout (b[0:N]) \
+  num_gangs (3 + a[3]) num_workers (5 + a[5]) vector_length (7 + a[7])
+  /* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */
   {
     for (COUNTERTYPE i = 0; i < N; i++)
       b[i] = i * 4;
   }
 
-#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N])
+  /* Parallelism dimensions: literal.  */
+#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) \
+  num_gangs (3) num_workers (5) vector_length (7)
+  /* { dg-prune-output "using vector_length \\(32\\), ignoring 7" } */
   {
     for (COUNTERTYPE ii = 0; ii < N; ii++)
       c[ii] = a[ii] + b[ii];
   }
 
   for (COUNTERTYPE i = 0; i < N; i++)
-    if (c[i] != a[i] + b[i])
-      abort ();
+    {
+      if (a[i] != i * 2)
+	abort ();
+      if (b[i] != i * 4)
+	abort ();
+      if (c[i] != a[i] + b[i])
+	abort ();
+    }
 
   free (a);
   free (b);
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
index d8af546..8308f7c 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-dims.c
@@ -520,5 +520,40 @@ int main ()
   }
 
 
+  /* Unparallelized OpenACC kernels constructs must get launched as 1 x 1 x 1
+     kernels even when there are explicit num_gangs, num_workers, or
+     vector_length clauses.  */
+  {
+    int gangs = 5;
+#define WORKERS 5
+#define VECTORS 13
+    int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
+    gangs_min = workers_min = vectors_min = INT_MAX;
+    gangs_max = workers_max = vectors_max = INT_MIN;
+#pragma acc kernels \
+  num_gangs (gangs) \
+  num_workers (WORKERS) \
+  vector_length (VECTORS)
+    {
+      /* This is to make the OpenACC kernels construct unparallelizable.  */
+      asm volatile ("" : : : "memory");
+
+#pragma acc loop reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
+      for (int i = 100; i > -100; --i)
+	{
+	  gangs_min = gangs_max = acc_gang ();
+	  workers_min = workers_max = acc_worker ();
+	  vectors_min = vectors_max = acc_vector ();
+	}
+    }
+    if (gangs_min != 0 || gangs_max != 1 - 1
+	|| workers_min != 0 || workers_max != 1 - 1
+	|| vectors_min != 0 || vectors_max != 1 - 1)
+      __builtin_abort ();
+#undef VECTORS
+#undef WORKERS
+  }
+
+
   return 0;
 }
diff --git libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95 libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
index 163e8d5..b88ca67 100644
--- libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
+++ libgomp/testsuite/libgomp.oacc-fortran/kernels-loop-2.f95
@@ -6,25 +6,34 @@ program main
   integer, dimension (0:n-1) :: a, b, c
   integer                    :: i, ii
 
+  ! Parallelism dimensions: compiler/runtime decides.
   !$acc kernels copyout (a(0:n-1))
   do i = 0, n - 1
      a(i) = i * 2
   end do
   !$acc end kernels
 
-  !$acc kernels copyout (b(0:n-1))
+  ! Parallelism dimensions: variable.
+  !$acc kernels copyout (b(0:n-1)) &
+  !$acc num_gangs (3 + a(3)) num_workers (5 + a(5)) vector_length (7 + a(7))
+  ! { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" }
   do i = 0, n -1
      b(i) = i * 4
   end do
   !$acc end kernels
 
-  !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))
+  ! Parallelism dimensions: literal.
+  !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1)) &
+  !$acc num_gangs (3) num_workers (5) vector_length (7)
+  ! { dg-prune-output "using vector_length \\(32\\), ignoring 7" }
   do ii = 0, n - 1
      c(ii) = a(ii) + b(ii)
   end do
   !$acc end kernels
 
   do i = 0, n - 1
+     if (a(i) .ne. i * 2) call abort
+     if (b(i) .ne. i * 4) call abort
      if (c(i) .ne. a(i) + b(i)) call abort
   end do
 


Grüße
 Thomas

^ permalink raw reply	[flat|nested] 5+ messages in thread

end of thread, other threads:[~2017-05-23 15:49 UTC | newest]

Thread overview: 5+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2017-05-11 12:28 OpenACC 2.5 kernels construct: num_gangs, num_workers, vector_length clauses Thomas Schwinge
2017-05-14 10:40 ` Thomas Schwinge
2017-05-19 11:09 ` Thomas Schwinge
2017-05-23  9:54 ` Jakub Jelinek
2017-05-23 15:58   ` Thomas Schwinge

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).