This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.

Index Nav:	[Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav:	[Date Prev] [Date Next]	[Thread Prev] [Thread Next]
Other format:	[Raw text]

[PATCH 10/10, OpenACC] Make new OpenACC kernels conversion the default; adjust and add tests

From: Kwok Cheung Yeung <kcy at codesourcery dot com>
To: <gcc-patches at gcc dot gnu dot org>, Jakub Jelinek <jakub at redhat dot com>
Cc: Thomas Schwinge <thomas at codesourcery dot com>
Date: Wed, 17 Jul 2019 22:32:06 +0100
Subject: [PATCH 10/10, OpenACC] Make new OpenACC kernels conversion the default; adjust and add tests
References: <5e191259-d5d5-34ce-7fd5-fc8d2e6d982e@codesourcery.com>

This patch makes the new kernel conversion scheme the default, and adjusts thetests accordingly.


2019-07-16  Thomas Schwinge  <thomas@codesourcery.com>
            Kwok Cheung Yeung  <kcy@codesourcery.com>

	gcc/c-family/
	* c.opt (fopenacc-kernels): Default to "split".

	gcc/fortran/
	* lang.opt (fopenacc-kernels): Default to "split".

	gcc/
	* doc/invoke.texi (-fopenacc-kernels): Update.

	gcc/testsuite/
	* c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c:
	New file.
	* c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c:
	Likewise.
	* c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c:
	Likewise.
	* c-c++-common/goacc/note-parallelism-1-kernels-loops.c: Likewise.
	* c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c:
	Likewise.
	* c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c:
	Likewise.
	* c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c:
	Likewise.
	* c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c:
	Likewise.
	* c-c++-common/goacc/note-parallelism-kernels-loop-auto.c:
	Likewise.
	* c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c:
	Likewise.
	* c-c++-common/goacc/note-parallelism-kernels-loops.c: Likewise.
	* c-c++-common/goacc/classify-kernels-unparallelized.c: Update.
	* c-c++-common/goacc/classify-kernels.c: Likewise.
	* c-c++-common/goacc/classify-parallel.c: Likewise.
	* c-c++-common/goacc/classify-routine.c: Likewise.
	* c-c++-common/goacc/if-clause-2.c: Likewise.
	* c-c++-common/goacc/kernels-conversion.c: Likewise.
	* c-c++-common/goacc/kernels-decompose-1.c: Likewise.
	* c-c++-common/goacc/loop-2-kernels.c: Likewise.
	* c-c++-common/goacc/note-parallelism.c: Likewise.
	* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
	* gfortran.dg/goacc/kernels-conversion.f95: Likewise.
	* gfortran.dg/goacc/kernels-decompose-1.f95: Likewise.
	* gfortran.dg/goacc/kernels-tree.f95: Likewise.
	* gfortran.dg/goacc/classify-kernels-unparallelized.f95
	* gfortran.dg/goacc/classify-kernels.f95
	* gfortran.dg/goacc/loop-2-kernels.f95

	libgomp/
	* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c:
	Update.
	* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
	Likewise.
---
 gcc/c-family/c.opt                                 |   2 +-
 gcc/doc/invoke.texi                                |   2 +-
 gcc/fortran/lang.opt                               |   2 +-
 .../goacc/classify-kernels-unparallelized.c        |   9 +-
 .../c-c++-common/goacc/classify-kernels.c          |   4 +-
 .../c-c++-common/goacc/classify-parallel.c         |   2 +-
 .../c-c++-common/goacc/classify-routine.c          |   2 +-
 gcc/testsuite/c-c++-common/goacc/if-clause-2.c     |   1 -
 .../c-c++-common/goacc/kernels-conversion.c        |  10 +-
 .../c-c++-common/goacc/kernels-decompose-1.c       |  69 ++++---
 gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c  |  14 +-
 ...sm-1-kernels-conditional-loop-independent_seq.c | 129 +++++++++++++
 .../goacc/note-parallelism-1-kernels-loop-auto.c   | 126 +++++++++++++
 ...te-parallelism-1-kernels-loop-independent_seq.c | 126 +++++++++++++
 .../goacc/note-parallelism-1-kernels-loops.c       |  47 +++++
 .../note-parallelism-1-kernels-straight-line.c     |  82 +++++++++
 .../note-parallelism-combined-kernels-loop-auto.c  | 121 ++++++++++++
 ...llelism-combined-kernels-loop-independent_seq.c | 121 ++++++++++++
 ...lism-kernels-conditional-loop-independent_seq.c | 204 +++++++++++++++++++++
 .../goacc/note-parallelism-kernels-loop-auto.c     | 138 ++++++++++++++
 ...note-parallelism-kernels-loop-independent_seq.c | 138 ++++++++++++++
 .../goacc/note-parallelism-kernels-loops.c         |  50 +++++
 .../c-c++-common/goacc/note-parallelism.c          |   3 +-
 .../c-c++-common/goacc/uninit-dim-clause.c         |   6 +-
 .../goacc/classify-kernels-unparallelized.f95      |   1 +
 .../gfortran.dg/goacc/classify-kernels.f95         |   1 +
 .../gfortran.dg/goacc/kernels-conversion.f95       |   7 +-
 .../gfortran.dg/goacc/kernels-decompose-1.f95      |  79 ++++----
 gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95   |   1 -
 gcc/testsuite/gfortran.dg/goacc/loop-2-kernels.f95 |  22 +--
 .../libgomp.oacc-c-c++-common/acc_prof-kernels-1.c |  17 +-
 .../kernels-decompose-1.c                          |   9 +-
 32 files changed, 1416 insertions(+), 129 deletions(-)

create mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.ccreate mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.ccreate mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.ccreate mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.ccreate mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.ccreate mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.ccreate mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.ccreate mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.ccreate mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.ccreate mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.ccreate mode 100644gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c


diff --git a/gcc/c-family/c.opt b/gcc/c-family/c.opt
index a193875..8efc5ea 100644
--- a/gcc/c-family/c.opt
+++ b/gcc/c-family/c.opt
@@ -1689,7 +1689,7 @@ C ObjC C++ ObjC++ LTO Joined Var(flag_openacc_dims)
 Specify default OpenACC compute dimensions.

 fopenacc-kernels=

-C ObjC C++ ObjC++ RejectNegative Joined Enum(openacc_kernels)Var(flag_openacc_kernels) Init(OPENACC_KERNELS_PARLOOPS)+C ObjC C++ ObjC++ RejectNegative Joined Enum(openacc_kernels)Var(flag_openacc_kernels) Init(OPENACC_KERNELS_SPLIT)-fopenacc-kernels=[split|parloops] Configure OpenACC 'kernels' constructshandling.


 Enum
diff --git a/gcc/doc/invoke.texi b/gcc/doc/invoke.texi
index ec98ab6..ffde9a2 100644
--- a/gcc/doc/invoke.texi
+++ b/gcc/doc/invoke.texi
@@ -2200,9 +2200,9 @@ Configure OpenACC 'kernels' constructs handling.
 With @option{-fopenacc-kernels=split}, OpenACC 'kernels' constructs
 are split into a sequence of compute constructs, each then handled
 individually.
+This is the default.
 With @option{-fopenacc-kernels=parloops}, the whole OpenACC
 'kernels' constructs is handled by the @samp{parloops} pass.
-This is the default.

 @item -fopenmp
 @opindex fopenmp
diff --git a/gcc/fortran/lang.opt b/gcc/fortran/lang.opt
index e7e277a..c84b284 100644
--- a/gcc/fortran/lang.opt
+++ b/gcc/fortran/lang.opt
@@ -663,7 +663,7 @@ Fortran LTO Joined Var(flag_openacc_dims)
 ; Documented in C

 fopenacc-kernels=

-Fortran RejectNegative Joined Enum(openacc_kernels) Var(flag_openacc_kernels)Init(OPENACC_KERNELS_PARLOOPS)+Fortran RejectNegative Joined Enum(openacc_kernels) Var(flag_openacc_kernels)Init(OPENACC_KERNELS_SPLIT)

 ; Documented in C

 fopenmp

diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.cb/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c

index d4c4b2c..443b207 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
@@ -1,5 +1,5 @@
 /* Check offloaded function's attributes and classification for unparallelized
-   OpenACC kernels.  */
+   OpenACC 'kernels'.  */

 /* { dg-additional-options "-O2" }
    { dg-additional-options "-fopt-info-optimized-omp" }
@@ -13,14 +13,15 @@ extern unsigned int *__restrict a;
 extern unsigned int *__restrict b;
 extern unsigned int *__restrict c;

-/* An "extern"al mapping of loop iterations/array indices makes the loop
-   unparallelizable.  */
 extern unsigned int f (unsigned int);
+#pragma acc routine (f) seq

 void KERNELS ()
 {

#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message"optimized: assigned OpenACC seq loop parallelism" } */

-  for (unsigned int i = 0; i < N; i++)

+ for (unsigned int i = 0; i < N; i++) /* { dg-message "optimized: beginning\"parloops\" region in OpenACC 'kernels' construct" } */

+    /* An "extern"al mapping of loop iterations/array indices makes the loop
+       unparallelizable.  */
     c[i] = a[f (i)] + b[f (i)];
 }

diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.cb/gcc/testsuite/c-c++-common/goacc/classify-kernels.c

index 16e9b9e..c154edf 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
@@ -1,5 +1,5 @@
 /* Check offloaded function's attributes and classification for OpenACC
-   kernels.  */
+   'kernels'.  */

 /* { dg-additional-options "-O2" }
    { dg-additional-options "-fopt-info-optimized-omp" }
@@ -16,7 +16,7 @@ extern unsigned int *__restrict c;
 void KERNELS ()
 {

#pragma acc kernels copyin (a[0:N], b[0:N]) copyout (c[0:N]) /* { dg-message"optimized: assigned OpenACC gang loop parallelism" } */

-  for (unsigned int i = 0; i < N; i++)

+ for (unsigned int i = 0; i < N; i++) /* { dg-message "optimized: beginning\"parloops\" region in OpenACC 'kernels' construct" } */

     c[i] = a[i] + b[i];
 }

diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.cb/gcc/testsuite/c-c++-common/goacc/classify-parallel.c

index 66a6d13..9c80efd 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c
@@ -1,5 +1,5 @@
 /* Check offloaded function's attributes and classification for OpenACC
-   parallel.  */
+   'parallel'.  */

 /* { dg-additional-options "-O2" }
    { dg-additional-options "-fopt-info-optimized-omp" }

diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.cb/gcc/testsuite/c-c++-common/goacc/classify-routine.c

index 0b9ba6e..a4994b0 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
@@ -1,5 +1,5 @@
 /* Check offloaded function's attributes and classification for OpenACC
-   routine.  */
+   'routine'.  */

 /* { dg-additional-options "-O2" }
    { dg-additional-options "-fopt-info-optimized-omp" }

diff --git a/gcc/testsuite/c-c++-common/goacc/if-clause-2.cb/gcc/testsuite/c-c++-common/goacc/if-clause-2.c

index e17b5dd..9920b4f 100644
--- a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c
@@ -1,4 +1,3 @@
-/* { dg-additional-options "-fopenacc-kernels=split" } */
 /* { dg-additional-options "-fdump-tree-convert_oacc_kernels" } */

 void

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-conversion.cb/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c

index ea7eec9..8cb63f0 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-conversion.c
@@ -1,4 +1,3 @@
-/* { dg-additional-options "-fopenacc-kernels=split" } */
 /* { dg-additional-options "-fdump-tree-convert_oacc_kernels" } */

 #define N 1024
@@ -52,12 +51,11 @@ main (void)
    parallelized loop region; and three "old-style" kernel regions. */

/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single" 1"convert_oacc_kernels" } } *//* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized" 1"convert_oacc_kernels" } } */

-/* { dg-final { scan-tree-dump-times "oacc_kernels" 3 "convert_oacc_kernels" } } */

+/* { dg-final { scan-tree-dump-times "oacc_kernels " 3 "convert_oacc_kernels" }} */


 /* Each of the parallel regions is async, and there is a final call to
    __builtin_GOACC_wait.  */

-/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels.* async$-1$" 5"convert_oacc_kernels" } } */+/* { dg-final { scan-tree-dump-times "oacc_kernels async\$-1\$" 3"convert_oacc_kernels" } } */+/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_singleasync\$-1\$" 1 "convert_oacc_kernels" } } */+/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelizedasync\$-1\$" 1 "convert_oacc_kernels" } } *//* { dg-final { scan-tree-dump-times "__builtin_GOACC_wait" 1"convert_oacc_kernels" } } */

-
-/* Check that the original kernels region is removed.  */
-/* { dg-final { scan-tree-dump-not "oacc_kernels" "convert_oacc_kernels" } } */

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.cb/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c

index b5d58c3..293ee42 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c
@@ -1,6 +1,5 @@
 /* Test OpenACC 'kernels' construct decomposition.  */

-/* { dg-additional-options "-fopenacc-kernels=split" } */
 /* { dg-additional-options "-fopt-info-optimized-omp" } */
 /* { dg-additional-options "-O2" } for "parloops".  */

@@ -31,92 +30,92 @@ main ()

 #pragma acc kernels
   {

- x = 0; /* { dg-message "note: beginning .gang-single. region in OpenACC.kernels. construct" } */+ x = 0; /* { dg-message "optimized: beginning .gang-single. region inOpenACC .kernels. construct" } */

     y = x < 10;
     z = x++;
     ;
   }

-#pragma acc kernels

- for (int i = 0; i < N; i++) /* { dg-message "note: beginning .parloops.region in OpenACC .kernels. construct" } */+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC gang loopparallelism" } */+ for (int i = 0; i < N; i++) /* { dg-message "optimized: beginning .parloops.region in OpenACC .kernels. construct" } */

     a[i] = 0;

-#pragma acc kernels loop

- /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to.parloops. for analysis" "" { target *-*-* } .-1 } */+#pragma acc kernels loop /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

   for (int i = 0; i < N; i++)
     b[i] = a[N - i - 1];

 #pragma acc kernels
   {
-#pragma acc loop

- /* { dg-message "note: forwarded loop nest in OpenACC .kernels. constructto .parloops. for analysis" "" { target *-*-* } .-1 } */+#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

     for (int i = 0; i < N; i++)
       b[i] = a[N - i - 1];

-#pragma acc loop

     for (int i = 0; i < N; i++)
       c[i] = a[i] * b[i];

- a[z] = 0; /* { dg-message "note: beginning .gang-single. region in OpenACC.kernels. construct" } */+ a[z] = 0; /* { dg-message "optimized: beginning .gang-single. region inOpenACC .kernels. construct" } */


-#pragma acc loop

     for (int i = 0; i < N; i++)
       c[i] += a[i];

-#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loopparallelism" } */- /* { dg-message "note: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

     for (int i = 0 + 1; i < N; i++)
       c[i] += c[i - 1];
   }

-#pragma acc kernels

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC worker vectorloop parallelism" } */

-#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang loopparallelism" } */- /* { dg-message "note: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangloop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

     for (int i = 0; i < N; ++i)

-#pragma acc loop independent /* { dg-message "note: assigned OpenACC workerloop parallelism" } */+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCworker loop parallelism" } */

       for (int j = 0; j < N; ++j)

-#pragma acc loop independent /* { dg-message "note: assigned OpenACC seq loopparallelism" } */+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC seqloop parallelism" } *//* { dg-warning "insufficient partitioning available to parallelize loop" ""{ target *-*-* } .-1 } */

 	for (int k = 0; k < N; ++k)
 	  a[(i + j + k) % N]
 	    = b[j]

- + f_v (c[k]); /* { dg-message "note: assigned OpenACC vector loopparallelism" } */+ + f_v (c[k]); /* { dg-message "optimized: assigned OpenACC vector loopparallelism" } */


     //TODO Should the following turn into "gang-single" instead of "parloops"?

//TODO The problem is that the first STMT is "if (y <= 4) goto <D.2547>;else goto <D.2548>;", thus "parloops".- if (y < 5) /* { dg-message "note: beginning .parloops. region in OpenACC.kernels. construct" } */-#pragma acc loop independent /* { dg-message "note: unparallelized loop nest inOpenACC .kernels. region: it's executed conditionally" } */+ if (y < 5) /* { dg-message "optimized: beginning .parloops. region inOpenACC .kernels. construct" } */+#pragma acc loop independent /* { dg-message "optimized: unparallelized loopnest in OpenACC .kernels. region: it's executed conditionally" } */

       for (int j = 0; j < N; ++j)
 	b[j] = f_w (c[j]);
   }

-#pragma acc kernels /* { dg-warning "region contains gang partitoned code butis not gang partitioned" } */

+#pragma acc kernels
   {

- /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels.construct" "" { target *-*-* } .+1 } */- y = f_g (a[5]); /* { dg-message "note: assigned OpenACC gang worker vectorloop parallelism" } */+ /* { dg-message "optimized: beginning .gang-single. region in OpenACC.kernels. construct" "" { target *-*-* } .+1 } */+ y = f_g (a[5]); /* { dg-message "optimized: assigned OpenACC gang workervector loop parallelism" } */

     for (int j = 0; j < N; ++j)

- b[j] = y + f_w (c[j]); /* { dg-message "note: assigned OpenACC workervector loop parallelism" } */+ b[j] = y + f_w (c[j]); /* { dg-message "optimized: assigned OpenACCworker vector loop parallelism" } */

   }

 #pragma acc kernels
   {

- y = 3; /* { dg-message "note: beginning .gang-single. region in OpenACC.kernels. construct" } */+ y = 3; /* { dg-message "optimized: beginning .gang-single. region inOpenACC .kernels. construct" } */

-#pragma acc loop independent /* { dg-message "note: assigned OpenACC gangworker loop parallelism" } */- /* { dg-message "note: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangworker loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

     for (int j = 0; j < N; ++j)

- b[j] = y + f_v (c[j]); /* { dg-message "note: assigned OpenACC vectorloop parallelism" } */+ b[j] = y + f_v (c[j]); /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

- z = 2; /* { dg-message "note: beginning .gang-single. region in OpenACC.kernels. construct" } */+ z = 2; /* { dg-message "optimized: beginning .gang-single. region inOpenACC .kernels. construct" } */

-#pragma acc kernels /* { dg-message "note: beginning .gang-single. region inOpenACC .kernels. construct" } */+#pragma acc kernels /* { dg-message "optimized: beginning .gang-single. regionin OpenACC .kernels. construct" } */

   ;

   return 0;

diff --git a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.cb/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c

index 0151508..c989222 100644
--- a/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/loop-2-kernels.c
@@ -37,7 +37,7 @@ void K(void)
 	for (j = 0; j < 10; j++)
 	  { }
       }
-#pragma acc loop seq gang // { dg-error "'seq' overrides" }
+#pragma acc loop seq gang // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } }
     for (i = 0; i < 10; i++)
       { }

@@ -63,7 +63,7 @@ void K(void)
 	for (j = 0; j < 10; j++)
 	  { }
       }
-#pragma acc loop seq worker // { dg-error "'seq' overrides" }

+#pragma acc loop seq worker // { dg-error "'seq' overrides" "TODO" { xfail*-*-* } }

     for (i = 0; i < 10; i++)
       { }
 #pragma acc loop gang worker
@@ -92,7 +92,7 @@ void K(void)
 	for (j = 1; j < 10; j++)
 	  { }
       }
-#pragma acc loop seq vector // { dg-error "'seq' overrides" }

+#pragma acc loop seq vector // { dg-error "'seq' overrides" "TODO" { xfail*-*-* } }

     for (i = 0; i < 10; i++)
       { }
 #pragma acc loop gang vector
@@ -105,7 +105,7 @@ void K(void)
 #pragma acc loop auto
     for (i = 0; i < 10; i++)
       { }
-#pragma acc loop seq auto // { dg-error "'seq' overrides" }
+#pragma acc loop seq auto // { dg-error "'seq' overrides" "TODO" { xfail *-*-* } }
     for (i = 0; i < 10; i++)
       { }
 #pragma acc loop gang auto // { dg-error "'auto' conflicts" }
@@ -147,7 +147,7 @@ void K(void)
 #pragma acc kernels loop worker(num:5)
   for (i = 0; i < 10; i++)
     { }
-#pragma acc kernels loop seq worker // { dg-error "'seq' overrides" }

+#pragma acc kernels loop seq worker // { dg-error "'seq' overrides" "TODO" {xfail *-*-* } }

   for (i = 0; i < 10; i++)
     { }
 #pragma acc kernels loop gang worker
@@ -163,7 +163,7 @@ void K(void)
 #pragma acc kernels loop vector(length:5)
   for (i = 0; i < 10; i++)
     { }
-#pragma acc kernels loop seq vector // { dg-error "'seq' overrides" }

+#pragma acc kernels loop seq vector // { dg-error "'seq' overrides" "TODO" {xfail *-*-* } }

   for (i = 0; i < 10; i++)
     { }
 #pragma acc kernels loop gang vector
@@ -176,7 +176,7 @@ void K(void)
 #pragma acc kernels loop auto
   for (i = 0; i < 10; i++)
     { }
-#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" }

+#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" "TODO" {xfail *-*-* } }

   for (i = 0; i < 10; i++)
     { }
 #pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" }

diff --gita/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c

new file mode 100644
index 0000000..c21273a
--- /dev/null

+++b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c

@@ -0,0 +1,129 @@
+/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels'
+   construct containing conditionally executed 'loop' constructs with
+   'independent' or 'seq' clauses.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+extern int c;
+
+int
+main ()
+{
+  int x, y, z;
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop seq

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc loop independent gang

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc loop independent worker

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc loop independent vector

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc loop independent gang vector

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc loop independent gang worker

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc loop independent worker vector

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc loop independent gang worker vector

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc loop independent gang

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent worker
+    for (y = 0; y < 10; y++)
+#pragma acc loop independent vector
+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc loop independent

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc loop independent

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent
+    for (y = 0; y < 10; y++)
+      ;
+
+#pragma acc loop independent

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent
+    for (y = 0; y < 10; y++)
+#pragma acc loop independent
+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc loop seq

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent
+    for (y = 0; y < 10; y++)
+#pragma acc loop independent
+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc loop independent

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop seq
+    for (y = 0; y < 10; y++)
+#pragma acc loop independent
+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc loop independent

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent
+    for (y = 0; y < 10; y++)
+#pragma acc loop seq
+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc loop seq

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent
+    for (y = 0; y < 10; y++)
+#pragma acc loop seq
+      for (z = 0; z < 10; z++)
+	;
+ }
+
+  return 0;
+}

diff --gita/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c

new file mode 100644
index 0000000..eedc472
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c
@@ -0,0 +1,126 @@
+/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels'
+   construct containing 'loop' constructs with explicit or implicit 'auto'
+   clause.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+int
+main ()
+{
+  int x, y, z;
+
+#pragma acc kernels
+ /* Strangely indented to keep this similar to other test cases.  */
+ {

+#pragma acc loop /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop gang /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop worker /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop vector /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop gang vector /* { dg-message "optimized: assigned OpenACC seqloop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop gang worker /* { dg-message "optimized: assigned OpenACC seqloop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop worker vector /* { dg-message "optimized: assigned OpenACC seqloop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop gang worker vector /* { dg-message "optimized: assignedOpenACC seq loop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop worker
+    for (y = 0; y < 10; y++)
+#pragma acc loop vector
+      for (z = 0; z < 10; z++)
+	;
+

+#pragma acc loop auto /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+      ;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop auto
+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop auto
+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop
+    for (y = 0; y < 10; y++)
+#pragma acc loop auto
+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop
+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop
+      for (z = 0; z < 10; z++)
+	;
+ }
+
+  return 0;
+}

diff --gita/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c

new file mode 100644
index 0000000..dad1bdb
--- /dev/null

+++b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c

@@ -0,0 +1,126 @@
+/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels'
+   construct containing 'loop' constructs with 'independent' or 'seq'
+   clauses.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+int
+main ()
+{
+  int x, y, z;
+
+#pragma acc kernels
+ /* Strangely indented to keep this similar to other test cases.  */
+ {

+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop independent gang /* { dg-message "optimized: assigned OpenACCgang loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop independent worker /* { dg-message "optimized: assignedOpenACC worker loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop independent vector /* { dg-message "optimized: assignedOpenACC vector loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop independent gang vector /* { dg-message "optimized: assignedOpenACC gang vector loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop independent gang worker /* { dg-message "optimized: assignedOpenACC gang worker loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop independent worker vector /* { dg-message "optimized: assignedOpenACC worker vector loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop independent gang worker vector /* { dg-message "optimized:assigned OpenACC gang worker vector loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent worker /* { dg-message "optimized: assignedOpenACC worker loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent vector /* { dg-message "optimized: assignedOpenACC vector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangvector loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangworker loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+    for (y = 0; y < 10; y++)
+      ;
+

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangloop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCworker loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangworker loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)

+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangvector loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+ }
+
+  return 0;
+}

diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c

new file mode 100644
index 0000000..336be88
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c
@@ -0,0 +1,47 @@
+/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels'
+   construct containing loops.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+int
+main ()
+{
+  int x, y, z;
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */
+ {

+ for (x = 0; x < 10; x++) /* { dg-message "optimized: beginning .parloops.region in OpenACC .kernels. construct" } */

+    ;
+
+  for (x = 0; x < 10; x++)
+    ;
+
+  for (x = 0; x < 10; x++)
+    for (y = 0; y < 10; y++)
+      for (z = 0; z < 10; z++)
+	;
+
+  for (x = 0; x < 10; x++)
+    ;
+
+  for (x = 0; x < 10; x++)
+    for (y = 0; y < 10; y++)
+      ;
+
+  for (x = 0; x < 10; x++)
+    for (y = 0; y < 10; y++)
+      for (z = 0; z < 10; z++)
+	;
+
+  for (x = 0; x < 10; x++)
+    for (y = 0; y < 10; y++)
+      for (z = 0; z < 10; z++)
+	;
+ }
+
+  return 0;
+}

diff --gita/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c

new file mode 100644
index 0000000..07a1e32
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c
@@ -0,0 +1,82 @@
+/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels'
+   construct containing straight-line code.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+#pragma acc routine gang
+extern int
+f_g (int);
+
+#pragma acc routine worker
+extern int
+f_w (int);
+
+#pragma acc routine vector
+extern int
+f_v (int);
+
+#pragma acc routine seq
+extern int
+f_s (int);
+
+int
+main ()
+{
+  int x, y, z;
+
+#pragma acc kernels
+  {

+ x = 0; /* { dg-message "optimized: beginning .gang-single. region inOpenACC .kernels. construct" } */

+    y = x < 10;
+    z = x++;
+    ;
+
+    y = 0;
+    z = y < 10;

+ x -= f_g (y++); /* { dg-message "optimized: assigned OpenACC gang workervector loop parallelism" } */

+    ;
+

+ x = f_w (0); /* { dg-message "optimized: assigned OpenACC worker vectorloop parallelism" } */+ z = f_v (x < 10); /* { dg-message "optimized: assigned OpenACC vector loopparallelism" } */+ y -= f_s (x++); /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+    ;
+
+    x = 0;
+    y = x < 10;
+    z = (x++);
+    y = 0;
+    x = y < 10;
+    z += (y++);
+    ;
+
+    x = 0;

+ y += f_s (x < 10); /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+    x++;
+    y = 0;

+ y += f_v (y < 10); /* { dg-message "optimized: assigned OpenACC vector loopparallelism" } */

+    y++;
+    z = 0;

+ y += f_w (z < 10); /* { dg-message "optimized: assigned OpenACC workervector loop parallelism" } */

+    z++;
+    ;
+
+    x = 0;

+ y *= f_g ( /* { dg-message "optimized: assigned OpenACC gang worker vectorloop parallelism" } */+ f_w (x < 10) /* { dg-message "optimized: assigned OpenACC worker vectorloop parallelism" } */+ + f_g (x < 10) /* { dg-message "optimized: assigned OpenACC gang workervector loop parallelism" } */

+	      );
+    x++;
+    y = 0;
+    y *= y < 10;
+    y++;
+    z = 0;
+    y *= z < 10;
+    z++;
+    ;
+  }
+
+  return 0;
+}

diff --gita/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c

new file mode 100644
index 0000000..2241901
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c
@@ -0,0 +1,121 @@
+/* Test the output of "-fopt-info-optimized-omp" for combined OpenACC 'kernels
+   loop' constructs with explicit or implicit 'auto' clause.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+int
+main ()
+{
+  int x, y, z;
+

+#pragma acc kernels loop /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop gang /* { dg-message "optimized: assigned OpenACC seqloop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop worker /* { dg-message "optimized: assigned OpenACCseq loop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop vector /* { dg-message "optimized: assigned OpenACCseq loop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop gang vector /* { dg-message "optimized: assignedOpenACC seq loop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop gang worker /* { dg-message "optimized: assignedOpenACC seq loop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop worker vector /* { dg-message "optimized: assignedOpenACC seq loop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop gang worker vector /* { dg-message "optimized:assigned OpenACC seq loop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop worker
+    for (y = 0; y < 10; y++)
+#pragma acc loop vector
+      for (z = 0; z < 10; z++)
+	;
+

+#pragma acc kernels loop auto /* { dg-message "optimized: assigned OpenACC seqloop parallelism" } */+ /* { dg-message "optimized: forwarded loop nest in OpenACC .kernels.construct to .parloops. for analysis" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+      ;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop auto
+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop auto
+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop
+    for (y = 0; y < 10; y++)
+#pragma acc loop auto
+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop
+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop
+      for (z = 0; z < 10; z++)
+	;
+
+  return 0;
+}

diff --gita/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c

new file mode 100644
index 0000000..b743636
--- /dev/null

+++b/gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c

@@ -0,0 +1,121 @@
+/* Test the output of "-fopt-info-optimized-omp" for combined OpenACC 'kernels
+   loop' constructs with 'independent' or 'seq' clauses.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+int
+main ()
+{
+  int x, y, z;
+

+#pragma acc kernels loop seq /* { dg-message "optimized: assigned OpenACC seqloop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop independent gang /* { dg-message "optimized: assignedOpenACC gang loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop independent worker /* { dg-message "optimized:assigned OpenACC worker loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop independent vector /* { dg-message "optimized:assigned OpenACC vector loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop independent gang vector /* { dg-message "optimized:assigned OpenACC gang vector loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop independent gang worker /* { dg-message "optimized:assigned OpenACC gang worker loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop independent worker vector /* { dg-message "optimized:assigned OpenACC worker vector loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop independent gang worker vector /* { dg-message"optimized: assigned OpenACC gang worker vector loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent worker /* { dg-message "optimized: assignedOpenACC worker loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent vector /* { dg-message "optimized: assignedOpenACC vector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+

+#pragma acc kernels loop independent /* { dg-message "optimized: assignedOpenACC gang vector loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+

+#pragma acc kernels loop independent /* { dg-message "optimized: assignedOpenACC gang worker loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+    for (y = 0; y < 10; y++)
+      ;
+

+#pragma acc kernels loop independent /* { dg-message "optimized: assignedOpenACC gang loop parallelism" } */+ /* { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCworker loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangworker loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)

+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangvector loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+
+  return 0;
+}

diff --gita/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c

new file mode 100644
index 0000000..21d31c4
--- /dev/null

+++b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c

@@ -0,0 +1,204 @@
+/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'kernels'
+   constructs containing conditionally executed 'loop' constructs with
+   'independent' or 'seq' clauses.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+extern int c;
+
+int
+main ()
+{
+  int x, y, z;
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop seq

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent gang

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent worker

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent vector

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent gang vector

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent gang worker

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent worker vector

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent gang worker vector

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent gang

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent worker
+    for (y = 0; y < 10; y++)
+#pragma acc loop independent vector
+      for (z = 0; z < 10; z++)
+	;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+    ;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent
+    for (y = 0; y < 10; y++)
+      ;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent
+    for (y = 0; y < 10; y++)
+#pragma acc loop independent
+      for (z = 0; z < 10; z++)
+	;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop seq

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent
+    for (y = 0; y < 10; y++)
+#pragma acc loop independent
+      for (z = 0; z < 10; z++)
+	;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop seq
+    for (y = 0; y < 10; y++)
+#pragma acc loop independent
+      for (z = 0; z < 10; z++)
+	;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop independent

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent
+    for (y = 0; y < 10; y++)
+#pragma acc loop seq
+      for (z = 0; z < 10; z++)
+	;
+ }
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+ /* Strangely indented to keep this similar to other test cases.  */

+ if (c) /* { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" } */

+ {
+#pragma acc loop seq

+ /* { dg-message "optimized: unparallelized loop nest in OpenACC .kernels.region: it's executed conditionally" "" { target *-*-* } .-1 } */

+  for (x = 0; x < 10; x++)
+#pragma acc loop independent
+    for (y = 0; y < 10; y++)
+#pragma acc loop seq
+      for (z = 0; z < 10; z++)
+	;
+ }
+
+  return 0;
+}

diff --gita/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c

new file mode 100644
index 0000000..02b9064
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c
@@ -0,0 +1,138 @@
+/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'kernels'
+   constructs containing 'loop' constructs with explicit or implicit 'auto'
+   clause.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+int
+main ()
+{
+  int x, y, z;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+#pragma acc loop worker
+    for (y = 0; y < 10; y++)
+#pragma acc loop vector
+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+      ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop auto
+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop auto
+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+#pragma acc loop
+    for (y = 0; y < 10; y++)
+#pragma acc loop auto
+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop
+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+#pragma acc loop auto
+    for (y = 0; y < 10; y++)
+#pragma acc loop
+      for (z = 0; z < 10; z++)
+	;
+
+  return 0;
+}

diff --gita/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c

new file mode 100644
index 0000000..6824d70
--- /dev/null

+++b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c

@@ -0,0 +1,138 @@
+/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'kernels'
+   constructs containing 'loop' constructs with 'independent' or 'seq'
+   clauses.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+int
+main ()
+{
+  int x, y, z;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent worker /* { dg-message "optimized: assignedOpenACC worker loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent vector /* { dg-message "optimized: assignedOpenACC vector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)
+    ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+    for (y = 0; y < 10; y++)
+      ;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCworker loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangworker loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)

+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACCvector loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+
+#pragma acc kernels

+  for (x = 0; x < 10; x++)

+#pragma acc loop independent /* { dg-message "optimized: assigned OpenACC gangvector loop parallelism" } */

+    for (y = 0; y < 10; y++)

+#pragma acc loop seq /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */

+      for (z = 0; z < 10; z++)
+	;
+
+  return 0;
+}

diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c

new file mode 100644
index 0000000..365464b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c
@@ -0,0 +1,50 @@
+/* Test the output of "-fopt-info-optimized-omp" for an OpenACC 'kernels'
+   construct containing loops.  */
+
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
+//TODO update accordingly
+/* See also "../../gfortran.dg/goacc/note-parallelism.f90".  */
+
+int
+main ()
+{
+  int x, y, z;
+

+#pragma acc kernels /* { dg-message "optimized: assigned OpenACC seq loopparallelism" } */+ for (x = 0; x < 10; x++) /* { dg-message "optimized: beginning .parloops.region in OpenACC .kernels. construct" } */

+    ;
+

+    ;
+

+    for (y = 0; y < 10; y++)
+      for (z = 0; z < 10; z++)
+	;
+

+    ;
+

+    for (y = 0; y < 10; y++)
+      ;
+

+    for (y = 0; y < 10; y++)
+      for (z = 0; z < 10; z++)
+	;
+

+    for (y = 0; y < 10; y++)
+      for (z = 0; z < 10; z++)
+	;
+
+  return 0;
+}

diff --git a/gcc/testsuite/c-c++-common/goacc/note-parallelism.cb/gcc/testsuite/c-c++-common/goacc/note-parallelism.c

index 735df7d..2b49a8b 100644
--- a/gcc/testsuite/c-c++-common/goacc/note-parallelism.c
+++ b/gcc/testsuite/c-c++-common/goacc/note-parallelism.c
@@ -1,4 +1,5 @@
-/* Test the output of "-fopt-info-optimized-omp".  */
+/* Test the output of "-fopt-info-optimized-omp" for OpenACC 'parallel'
+   constructs.  */

 /* { dg-additional-options "-fopt-info-optimized-omp" } */

diff --git a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.cb/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c

index 9f11196..f00daa7 100644
--- a/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
+++ b/gcc/testsuite/c-c++-common/goacc/uninit-dim-clause.c
@@ -18,12 +18,12 @@ void acc_kernels()
 {
   int i, j, k;

- #pragma acc kernels num_gangs(i) /* { dg-warning "is used uninitialized inthis function" } */+ #pragma acc kernels num_gangs(i) /* { dg-warning "is used uninitialized inthis function" "TODO" { xfail *-*-* } } */

- #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized inthis function" } */+ #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized inthis function" "TODO" { xfail *-*-* } } */

- #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitializedin this function" } */+ #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitializedin this function" "TODO" { xfail *-*-* } } */

   ;
 }

diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95

index 0877242..27ba39b 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels-unparallelized.f95
@@ -21,6 +21,7 @@ program main

   !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))

do i = 0, n - 1 ! { dg-message "optimized: assigned OpenACC seq loopparallelism" }+ ! { dg-message "optimized: beginning \"parloops\" region in OpenACC 'kernels'construct" "" { target *-*-* } .-1 }

      c(i) = a(f (i)) + b(f (i))
   end do
   !$acc end kernels

diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95

index f2c4736..68d0512 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-kernels.f95
@@ -17,6 +17,7 @@ program main

   !$acc kernels copyin (a(0:n-1), b(0:n-1)) copyout (c(0:n-1))

do i = 0, n - 1 ! { dg-message "optimized: assigned OpenACC gang loopparallelism" }+ ! { dg-message "optimized: beginning \"parloops\" region in OpenACC 'kernels'construct" "" { target *-*-* } .-1 }

      c(i) = a(i) + b(i)
   end do
   !$acc end kernels

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95b/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95

index 6604727..4672d15 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-conversion.f95
@@ -1,4 +1,3 @@
-! { dg-additional-options "-fopenacc-kernels=split" }
 ! { dg-additional-options "-fdump-tree-convert_oacc_kernels" }

 program main
@@ -50,9 +49,11 @@ end program main
 ! parallelized loop region; and three "old-style" kernel regions.

! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single" 1"convert_oacc_kernels" } }! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized" 1"convert_oacc_kernels" } }

-! { dg-final { scan-tree-dump-times "oacc_kernels" 3 "convert_oacc_kernels" } }
+! { dg-final { scan-tree-dump-times "oacc_kernels " 3 "convert_oacc_kernels" } }

 ! Each of the parallel regions is async, and there is a final call to
 ! __builtin_GOACC_wait.

-! { dg-final { scan-tree-dump-times "oacc_parallel_kernels.* async$-1$" 5"convert_oacc_kernels" } }+! { dg-final { scan-tree-dump-times "oacc_kernels async\$-1\$" 3"convert_oacc_kernels" } }+! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_singleasync\$-1\$" 1 "convert_oacc_kernels" } }+! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelizedasync\$-1\$" 1 "convert_oacc_kernels" } }! { dg-final { scan-tree-dump-times "__builtin_GOACC_wait" 1"convert_oacc_kernels" } }diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95

index 520bf03..b2956d7 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-decompose-1.f95
@@ -1,6 +1,5 @@
 ! Test OpenACC 'kernels' construct decomposition.

-! { dg-additional-options "-fopenacc-kernels=split" }
 ! { dg-additional-options "-fopt-info-optimized-omp" }
 ! { dg-additional-options "-O2" } for "parloops".

@@ -25,7 +24,7 @@ program main
   integer :: a(N), b(N), c(N)

   !$acc kernels

- x = 0 ! { dg-message "note: beginning .gang-single. region in OpenACC.kernels. construct" }+ x = 0 ! { dg-message "optimized: beginning .gang-single. region in OpenACC.kernels. construct" }

   y = 0
   y_l = x < 10
   z = x
@@ -33,67 +32,67 @@ program main
   ;
   !$acc end kernels

-  !$acc kernels ! { dg-message "note: assigned OpenACC gang loop parallelism" }

- do i = 1, N ! { dg-message "note: beginning .parloops. region in OpenACC.kernels. construct" }+ !$acc kernels ! { dg-message "optimized: assigned OpenACC gang loopparallelism" }+ do i = 1, N ! { dg-message "optimized: beginning .parloops. region in OpenACC.kernels. construct" }

      a(i) = 0
   end do
   !$acc end kernels

- !$acc kernels loop ! { dg-message "note: assigned OpenACC gang loopparallelism" }- ! { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to.parloops. for analysis" "" { target *-*-* } .-1 }+ !$acc kernels loop ! { dg-message "optimized: assigned OpenACC seq loopparallelism" }+ ! { dg-message "optimized: forwarded loop nest in OpenACC .kernels. constructto .parloops. for analysis" "" { target *-*-* } .-1 }

   do i = 1, N
      b(i) = a(N - i + 1)
   end do

   !$acc kernels
-  !$acc loop ! { dg-message "note: assigned OpenACC gang loop parallelism" }

- ! { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to.parloops. for analysis" "" { target *-*-* } .-1 }

+  !$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }

+ ! { dg-message "optimized: forwarded loop nest in OpenACC .kernels. constructto .parloops. for analysis" "" { target *-*-* } .-1 }

   do i = 1, N
      b(i) = a(N - i + 1)
   end do

-  !$acc loop ! { dg-message "note: assigned OpenACC gang loop parallelism" }

- ! { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to.parloops. for analysis" "" { target *-*-* } .-1 }

+  !$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }

+ ! { dg-message "optimized: forwarded loop nest in OpenACC .kernels. constructto .parloops. for analysis" "" { target *-*-* } .-1 }

   do i = 1, N
      c(i) = a(i) * b(i)
   end do

- a(z) = 0 ! { dg-message "note: beginning .gang-single. region in OpenACC.kernels. construct" }+ a(z) = 0 ! { dg-message "optimized: beginning .gang-single. region in OpenACC.kernels. construct" }


-  !$acc loop ! { dg-message "note: assigned OpenACC gang loop parallelism" }

- ! { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to.parloops. for analysis" "" { target *-*-* } .-1 }

+  !$acc loop ! { dg-message "optimized: assigned OpenACC seq loop parallelism" }

+ ! { dg-message "optimized: forwarded loop nest in OpenACC .kernels. constructto .parloops. for analysis" "" { target *-*-* } .-1 }

   do i = 1, N
      c(i) = c(i) + a(i)
   end do

-  !$acc loop seq ! { dg-message "note: assigned OpenACC seq loop parallelism" }

- ! { dg-message "note: parallelized loop nest in OpenACC .kernels. construct""" { target *-*-* } .-1 }+ !$acc loop seq ! { dg-message "optimized: assigned OpenACC seq loopparallelism" }+ ! { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 }

   do i = 1 + 1, N
      c(i) = c(i) + c(i - 1)
   end do
   !$acc end kernels

- !$acc kernels ! { dg-bogus "note: assigned OpenACC seq loop parallelism""TODO" { xfail *-*-* } }- !$acc loop independent ! { dg-message "note: assigned OpenACC gang loopparallelism" }- ! { dg-message "note: parallelized loop nest in OpenACC .kernels. construct""" { target *-*-* } .-1 }+ !$acc kernels ! { dg-message "optimized: assigned OpenACC worker vector loopparallelism" }+ !$acc loop independent ! { dg-message "optimized: assigned OpenACC gang loopparallelism" }+ ! { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 }

   do i = 1, N

- !$acc loop independent ! { dg-message "note: assigned OpenACC worker loopparallelism" }+ !$acc loop independent ! { dg-message "optimized: assigned OpenACC workerloop parallelism" }

      do j = 1, N

- !$acc loop independent ! { dg-message "note: assigned OpenACC seq loopparallelism" "TODO" { xfail *-*-* } }- ! { dg-warning "insufficient partitioning available to parallelizeloop" "TODO" { xfail *-*-* } .-1 }- ! { dg-bogus "note: assigned OpenACC vector loop parallelism" "TODO" {xfail *-*-* } .-2 }+ !$acc loop independent ! { dg-message "optimized: assigned OpenACC seqloop parallelism" }+ ! { dg-warning "insufficient partitioning available to parallelizeloop" "" { target *-*-* } .-1 }

+        ! { dg-bogus "optimized: assigned OpenACC vector loop parallelism" }
         do k = 1, N
            a(1 + mod(i + j + k, N)) &
                 = b(j) &

- + f_v (c(k)) ! { dg-message "note: assigned OpenACC vector loopparallelism" "TODO" { xfail *-*-* } .-1 }+ + f_v (c(k)) ! { dg-message "optimized: assigned OpenACC vectorloop parallelism" }

         end do
      end do
   end do

   !TODO Should the following turn into "gang-single" instead of "parloops"?

!TODO The problem is that the first STMT is "if (y <= 4) goto <D.2547>; elsegoto <D.2548>;", thus "parloops".- if (y < 5) then ! { dg-message "note: beginning .parloops. region in OpenACC.kernels. construct" }- !$acc loop independent ! { dg-message "note: unparallelized loop nest inOpenACC .kernels. region: it's executed conditionally" }+ if (y < 5) then ! { dg-message "optimized: beginning .parloops. region inOpenACC .kernels. construct" }+ !$acc loop independent ! { dg-message "optimized: unparallelized loop nestin OpenACC .kernels. region: it's executed conditionally" }

      do j = 1, N
         b(j) = f_w (c(j))
      end do
@@ -101,32 +100,30 @@ program main
   !$acc end kernels

   !$acc kernels
-  !TODO This refers to the "gang-single" "f_g" call.

- ! { dg-warning "region contains gang partitoned code but is not gangpartitioned" "TODO" { xfail *-*-* } .-2 }- ! { dg-message "note: beginning .gang-single. region in OpenACC .kernels.construct" "" { target *-*-* } .+1 }- y = f_g (a(5)) ! { dg-message "note: assigned OpenACC gang worker vector loopparallelism" "TODO" { xfail *-*-* } }

- !$acc loop independent ! { dg-message "note: assigned OpenACC gang loopparallelism" "TODO" { xfail *-*-* } }- ! { dg-message "note: parallelized loop nest in OpenACC .kernels. construct""" { target *-*-* } .-1 }- ! { dg-bogus "note: assigned OpenACC gang vector loop parallelism" "TODO" {xfail *-*-* } .-2 }+ ! { dg-message "optimized: beginning .gang-single. region in OpenACC.kernels. construct" "" { target *-*-* } .+1 }+ y = f_g (a(5)) ! { dg-message "optimized: assigned OpenACC gang worker vectorloop parallelism" }

+ !$acc loop independent ! { dg-message "optimized: assigned OpenACC gang loopparallelism" }+ ! { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 }+ ! { dg-bogus "optimized: assigned OpenACC gang vector loop parallelism" "" {target *-*-* } .-2 }

   do j = 1, N

- b(j) = y + f_w (c(j)) ! { dg-message "note: assigned OpenACC worker vectorloop parallelism" "TODO" { xfail *-*-* } }+ b(j) = y + f_w (c(j)) ! { dg-message "optimized: assigned OpenACC workervector loop parallelism" }

   end do
   !$acc end kernels

   !$acc kernels

- y = 3 ! { dg-message "note: beginning .gang-single. region in OpenACC.kernels. construct" }+ y = 3 ! { dg-message "optimized: beginning .gang-single. region in OpenACC.kernels. construct" }

- !$acc loop independent ! { dg-message "note: assigned OpenACC gang workerloop parallelism" "TODO" { xfail *-*-* } }- ! { dg-message "note: parallelized loop nest in OpenACC .kernels. construct""" { target *-*-* } .-1 }- ! { dg-bogus "note: assigned OpenACC gang vector loop parallelism" "TODO" {xfail *-*-* } .-2 }+ !$acc loop independent ! { dg-message "optimized: assigned OpenACC gangworker loop parallelism" }+ ! { dg-message "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } .-1 }+ ! { dg-bogus "optimized: assigned OpenACC gang vector loop parallelism" "" {target *-*-* } .-2 }

   do j = 1, N

- b(j) = y + f_v (c(j)) ! { dg-message "note: assigned OpenACC vector loopparallelism" "TODO" { xfail *-*-* } }+ b(j) = y + f_v (c(j)) ! { dg-message "optimized: assigned OpenACC vectorloop parallelism" }

   end do

- z = 2 ! { dg-message "note: beginning .gang-single. region in OpenACC.kernels. construct" }+ z = 2 ! { dg-message "optimized: beginning .gang-single. region in OpenACC.kernels. construct" }

   !$acc end kernels

- !$acc kernels ! { dg-message "note: beginning .gang-single. region in OpenACC.kernels. construct" }+ !$acc kernels ! { dg-message "optimized: beginning .gang-single. region inOpenACC .kernels. construct" }

   !$acc end kernels
 end program main

diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95

index b83ca2d..bc9beba 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -1,6 +1,5 @@
 ! { dg-do compile }
 ! { dg-additional-options "-fdump-tree-original" }
-! { dg-additional-options "-fopenacc-kernels=split" }
 ! { dg-additional-options "-fdump-tree-convert_oacc_kernels" }

 program test

diff --git a/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels.f95b/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels.f95

index 874c62d..f77bb23 100644
--- a/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/loop-2-kernels.f95
@@ -35,7 +35,7 @@ program test
       DO j = 1,10
       ENDDO
     ENDDO

- !$acc loop seq gang ! { dg-error "'seq' overrides other OpenACC loopspecifiers" }+ !$acc loop seq gang ! { dg-error "'seq' overrides other OpenACC loopspecifiers" "TODO" { xfail *-*-* } }

     DO i = 1,10
     ENDDO

@@ -56,11 +56,11 @@ program test

!$acc loop worker ! { dg-error "inner loop uses same OpenACC parallelismas containing loop" }

       DO j = 1,10
       ENDDO
-      !$acc loop gang ! { dg-error "" "TODO" { xfail *-*-* } }
+      !$acc loop gang
       DO j = 1,10
       ENDDO
     ENDDO

- !$acc loop seq worker ! { dg-error "'seq' overrides other OpenACC loopspecifiers" }+ !$acc loop seq worker ! { dg-error "'seq' overrides other OpenACC loopspecifiers" "TODO" { xfail *-*-* } }

     DO i = 1,10
     ENDDO
     !$acc loop gang worker
@@ -81,14 +81,14 @@ program test

!$acc loop vector ! { dg-error "inner loop uses same OpenACC parallelismas containing loop" }

       DO j = 1,10
       ENDDO
-      !$acc loop worker ! { dg-error "" "TODO" { xfail *-*-* } }
+      !$acc loop worker
       DO j = 1,10
       ENDDO
-      !$acc loop gang ! { dg-error "" "TODO" { xfail *-*-* } }
+      !$acc loop gang
       DO j = 1,10
       ENDDO
     ENDDO

- !$acc loop seq vector ! { dg-error "'seq' overrides other OpenACC loopspecifiers" }+ !$acc loop seq vector ! { dg-error "'seq' overrides other OpenACC loopspecifiers" "TODO" { xfail *-*-* } }

     DO i = 1,10
     ENDDO
     !$acc loop gang vector
@@ -101,7 +101,7 @@ program test
     !$acc loop auto
     DO i = 1,10
     ENDDO

- !$acc loop seq auto ! { dg-error "'seq' overrides other OpenACC loopspecifiers" }+ !$acc loop seq auto ! { dg-error "'seq' overrides other OpenACC loopspecifiers" "TODO" { xfail *-*-* } }

     DO i = 1,10
     ENDDO

!$acc loop gang auto ! { dg-error "'auto' conflicts with other OpenACCloop specifiers" }

@@ -133,7 +133,7 @@ program test
   !$acc kernels loop gang(static:*)
   DO i = 1,10
   ENDDO

- !$acc kernels loop seq gang ! { dg-error "'seq' overrides other OpenACC loopspecifiers" }+ !$acc kernels loop seq gang ! { dg-error "'seq' overrides other OpenACC loopspecifiers" "TODO" { xfail *-*-* } }

   DO i = 1,10
   ENDDO

@@ -146,7 +146,7 @@ program test
   !$acc kernels loop worker(num:5)
   DO i = 1,10
   ENDDO

- !$acc kernels loop seq worker ! { dg-error "'seq' overrides other OpenACCloop specifiers" }+ !$acc kernels loop seq worker ! { dg-error "'seq' overrides other OpenACCloop specifiers" "TODO" { xfail *-*-* } }

   DO i = 1,10
   ENDDO
   !$acc kernels loop gang worker
@@ -162,7 +162,7 @@ program test
   !$acc kernels loop vector(length:5)
   DO i = 1,10
   ENDDO

- !$acc kernels loop seq vector ! { dg-error "'seq' overrides other OpenACCloop specifiers" }+ !$acc kernels loop seq vector ! { dg-error "'seq' overrides other OpenACCloop specifiers" "TODO" { xfail *-*-* } }

   DO i = 1,10
   ENDDO
   !$acc kernels loop gang vector
@@ -175,7 +175,7 @@ program test
   !$acc kernels loop auto
   DO i = 1,10
   ENDDO

- !$acc kernels loop seq auto ! { dg-error "'seq' overrides other OpenACC loopspecifiers" }+ !$acc kernels loop seq auto ! { dg-error "'seq' overrides other OpenACC loopspecifiers" "TODO" { xfail *-*-* } }

   DO i = 1,10
   ENDDO

!$acc kernels loop gang auto ! { dg-error "'auto' conflicts with otherOpenACC loop specifiers" }diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.cb/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c

index 7cfc364..fd339d2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c
@@ -41,6 +41,7 @@ static int state = -1;
 static acc_device_t acc_device_type;
 static int acc_device_num;
 static int num_gangs, num_workers, vector_length;
+static int async;

static void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info*event_info, acc_api_info *api_info)@@ -58,7 +59,7 @@ static void cb_enqueue_launch_start (acc_prof_info *prof_info,acc_event_info *e

   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 == async);
   assert (prof_info->async_queue == prof_info->async);
   assert (prof_info->src_file == NULL);
   assert (prof_info->func_name == NULL);
@@ -154,8 +155,10 @@ int main()
   acc_device_num = acc_get_device_num (acc_device_type);
   assert (state == 0);

-  /* Parallelism dimensions: compiler/runtime decides.  */
   STATE_OP (state, = 0);
+  /* Implicit async.  */
+  async = acc_async_noval;
+  /* Parallelism dimensions: compiler/runtime decides.  */
   num_gangs = num_workers = vector_length = 0;
   {
 #define N 100
@@ -175,8 +178,10 @@ int main()
 #undef N
   }

-  /* Parallelism dimensions: literal.  */
   STATE_OP (state, = 0);
+  /* Explicit async: without argument.  */
+  async = acc_async_noval;
+  /* Parallelism dimensions: literal.  */
   num_gangs = 30;
   num_workers = 3;
   vector_length = 5;
@@ -184,6 +189,7 @@ int main()
 #define N 100
     int x[N];
 #pragma acc kernels \
+  async \
   num_gangs (30) num_workers (3) vector_length (5)
     /* { dg-prune-output "using vector_length \\(32\\), ignoring 5" } */
     {
@@ -200,8 +206,10 @@ int main()
 #undef N
   }

-  /* Parallelism dimensions: variable.  */
   STATE_OP (state, = 0);
+  /* Explicit async: variable.  */
+  async = 123;
+  /* Parallelism dimensions: variable.  */
   num_gangs = 22;
   num_workers = 5;
   vector_length = 7;
@@ -209,6 +217,7 @@ int main()
 #define N 100
     int x[N];
 #pragma acc kernels \
+  async (async) \
   num_gangs (num_gangs) num_workers (num_workers) vector_length (vector_length)

/* { dg-prune-output "using vector_length \$32\$, ignoring runtimesetting" } */

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.cb/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c

index 601e543..45d786d 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c
@@ -1,4 +1,3 @@
-/* { dg-additional-options "-fopenacc-kernels=split" } */
 /* { dg-additional-options "-fopt-info-optimized-omp" } */

 #undef NDEBUG
@@ -12,14 +11,14 @@ int main()

 #pragma acc kernels
   {

- int c = 234; /* { dg-warning "note: beginning .gang-single. region inOpenACC .kernels. construct" } */+ int c = 234; /* { dg-warning "optimized: beginning .gang-single. region inOpenACC .kernels. construct" } */

-#pragma acc loop independent gang /* { dg-warning "note: assigned OpenACC gangloop parallelism" } */- /* { dg-warning "note: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } 17 } */+#pragma acc loop independent gang /* { dg-warning "optimized: assigned OpenACCgang loop parallelism" } */+ /* { dg-warning "optimized: parallelized loop nest in OpenACC .kernels.construct" "" { target *-*-* } 16 } */

     for (int i = 0; i < N; ++i)
       b[i] = c;

- a = c; /* { dg-warning "note: beginning .gang-single. region in OpenACC.kernels. construct" } */+ a = c; /* { dg-warning "optimized: beginning .gang-single. region inOpenACC .kernels. construct" } */

   }

   for (int i = 0; i < N; ++i)
--
2.8.1

References:
- [PATCH 00/10, OpenACC] Rework handling of OpenACC kernels regions
  - From: Kwok Cheung Yeung

Index Nav:	[Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav:	[Date Prev] [Date Next]	[Thread Prev] [Thread Next]