From 53f61640a510d13537709239d523c283881f0755 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed, 23 Jan 2019 02:40:08 -0800
Subject: [PATCH 9/9] Make new OpenACC kernels conversion the default; adjust
and add tests
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/dtype-1.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/routine-1.c: Likewise.
* c-c++-common/goacc/uninit-dim-clause.c: Likewise.
* gfortran.dg/goacc/dtype-1.f95: Likewise.
* gfortran.dg/goacc/kernels-conversion.f95: Likewise.
* gfortran.dg/goacc/kernels-decompose-1.f95: Likewise.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c:
Update.
* testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
Likewise.
* testsuite/libgomp.oacc-fortran/avoid-offloading-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/avoid-offloading-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/avoid-offloading-3.f: Likewise.
* testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90:
Likewise.
---
gcc/ChangeLog.openacc | 4 +
gcc/c-family/ChangeLog.openacc | 4 +
gcc/c-family/c.opt | 2 +-
gcc/doc/invoke.texi | 2 +-
gcc/fortran/ChangeLog.openacc | 4 +
gcc/fortran/lang.opt | 2 +-
gcc/testsuite/ChangeLog.openacc | 40 ++++
.../goacc/classify-kernels-unparallelized.c | 7 +-
.../c-c++-common/goacc/classify-kernels.c | 2 +-
.../c-c++-common/goacc/classify-parallel.c | 2 +-
.../c-c++-common/goacc/classify-routine.c | 2 +-
gcc/testsuite/c-c++-common/goacc/dtype-1.c | 8 +-
.../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 | 1 -
.../c-c++-common/goacc/loop-2-kernels.c | 14 +-
...kernels-conditional-loop-independent_seq.c | 129 +++++++++++
.../note-parallelism-1-kernels-loop-auto.c | 126 +++++++++++
...rallelism-1-kernels-loop-independent_seq.c | 126 +++++++++++
.../goacc/note-parallelism-1-kernels-loops.c | 47 ++++
...note-parallelism-1-kernels-straight-line.c | 82 +++++++
...e-parallelism-combined-kernels-loop-auto.c | 121 +++++++++++
...sm-combined-kernels-loop-independent_seq.c | 121 +++++++++++
...kernels-conditional-loop-independent_seq.c | 204 ++++++++++++++++++
.../note-parallelism-kernels-loop-auto.c | 138 ++++++++++++
...parallelism-kernels-loop-independent_seq.c | 138 ++++++++++++
.../goacc/note-parallelism-kernels-loops.c | 50 +++++
.../c-c++-common/goacc/note-parallelism.c | 3 +-
gcc/testsuite/c-c++-common/goacc/routine-1.c | 2 +-
.../c-c++-common/goacc/uninit-dim-clause.c | 6 +-
gcc/testsuite/gfortran.dg/goacc/dtype-1.f95 | 8 +-
.../gfortran.dg/goacc/kernels-conversion.f95 | 7 +-
.../gfortran.dg/goacc/kernels-decompose-1.f95 | 1 -
.../gfortran.dg/goacc/kernels-tree.f95 | 1 -
libgomp/ChangeLog.openacc | 16 ++
.../acc_prof-kernels-1.c | 17 +-
.../avoid-offloading-1.c | 18 +-
.../avoid-offloading-2.c | 17 +-
.../avoid-offloading-3.c | 14 +-
.../kernels-decompose-1.c | 3 +-
.../libgomp.oacc-fortran/avoid-offloading-1.f | 18 +-
.../libgomp.oacc-fortran/avoid-offloading-2.f | 18 +-
.../libgomp.oacc-fortran/avoid-offloading-3.f | 15 +-
.../initialize_kernels_loops.f90 | 1 -
44 files changed, 1494 insertions(+), 58 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-conditional-loop-independent_seq.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-auto.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loop-independent_seq.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-loops.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-1-kernels-straight-line.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-auto.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-combined-kernels-loop-independent_seq.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-conditional-loop-independent_seq.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-auto.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loop-independent_seq.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/note-parallelism-kernels-loops.c
@@ -1,3 +1,7 @@
+2019-01-31 Thomas Schwinge <thomas@codesourcery.com>
+
+ * doc/invoke.texi (-fopenacc-kernels): Update.
+
2019-01-31 Thomas Schwinge <thomas@codesourcery.com>
Gergö Barany <gergo@codesourcery.com>
@@ -1,3 +1,7 @@
+2019-01-31 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c.opt (fopenacc-kernels): Default to "split".
+
2019-01-31 Thomas Schwinge <thomas@codesourcery.com>
* c.opt (fopenacc-kernels): New flag.
@@ -1618,7 +1618,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' constructs handling.
Enum
@@ -2165,9 +2165,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
@@ -1,3 +1,7 @@
+2019-01-31 Thomas Schwinge <thomas@codesourcery.com>
+
+ * lang.opt (fopenacc-kernels): Default to "split".
+
2019-01-31 Thomas Schwinge <thomas@codesourcery.com>
* lang.opt (fopenacc-kernels): New flag.
@@ -643,7 +643,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
@@ -1,3 +1,43 @@
+2019-01-31 Thomas Schwinge <thomas@codesourcery.com>
+ Gergö Barany <gergo@codesourcery.com>
+
+ * 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/dtype-1.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/routine-1.c: Likewise.
+ * c-c++-common/goacc/uninit-dim-clause.c: Likewise.
+ * gfortran.dg/goacc/dtype-1.f95: Likewise.
+ * gfortran.dg/goacc/kernels-conversion.f95: Likewise.
+ * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise.
+ * gfortran.dg/goacc/kernels-tree.f95: Likewise.
+
2019-01-31 Thomas Schwinge <thomas@codesourcery.com>
Gergö Barany <gergo@codesourcery.com>
@@ -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 "note: assigned OpenACC seq loop parallelism" } */
for (unsigned int i = 0; i < N; i++)
+ /* An "extern"al mapping of loop iterations/array indices makes the loop
+ unparallelizable. */
c[i] = a[f (i)] + b[f (i)];
}
@@ -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" }
@@ -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" }
@@ -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" }
@@ -96,11 +96,13 @@ test ()
/* { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(\\*\\) \\\[ wait\\(10\\) vector_length\\(10\\) num_workers\\(10\\) num_gangs\\(10\\) async\\(10\\) \\\] device_type\\(nvidia\\) \\\[ wait\\(3\\) vector_length\\(128\\) num_workers\\(300\\) num_gangs\\(300\\) async\\(3\\) \\\] wait\\(1\\) vector_length\\(1\\) num_workers\\(1\\) num_gangs\\(1\\) async\\(1\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ wait\\(-1\\) async\\(-1\\) \\\]" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single async\\(-1\\) num_gangs\\(1\\) device_type\\(nvidia\\) \\\[ wait\\(-1\\) async\\(-1\\) \\\]" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ wait\\(1\\) async\\(1\\) \\\] wait\\(-1\\) async\\(-1\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(nvidia\\) \\\[ wait\\(1\\) async\\(1\\) \\\] wait\\(-1\\) async\\(-1\\)" 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(\\*\\) \\\[ wait\\(0\\) async\\(0\\) \\\] device_type\\(nvidia\\) \\\[ wait\\(2\\) async\\(2\\) \\\] wait\\(-1\\) async\\(-1\\)" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(\\*\\) \\\[ wait\\(0\\) async\\(0\\) \\\] device_type\\(nvidia\\) \\\[ wait\\(2\\) async\\(2\\) \\\] wait\\(-1\\) async\\(-1\\)" 1 "omplower" } } */
+
+/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(\\*\\) \\\[ wait\\(0\\) async\\(0\\) \\\] device_type\\(nvidia_ptx\\) \\\[ wait\\(1\\) async\\(1\\) \\\] wait\\(-1\\) async\\(-1\\)" 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ tile\\(1\\) gang \\\] private\\(i1\\.0\\) private\\(i1\\)" 1 "omplower" } } */
@@ -1,4 +1,3 @@
-/* { dg-additional-options "-fopenacc-kernels=split" } */
/* { dg-additional-options "-fdump-tree-convert_oacc_kernels" } */
void
@@ -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_single async\\(-1\\)" 1 "convert_oacc_kernels" } } */
+/* { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized async\\(-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" } } */
@@ -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". */
@@ -35,7 +35,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++)
{ }
@@ -61,7 +61,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
@@ -90,7 +90,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
@@ -103,7 +103,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
@@ -145,7 +145,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
@@ -161,7 +161,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
@@ -174,7 +174,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
new file mode 100644
@@ -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 "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop seq
+ /* { dg-message "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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 "note: 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;
+}
new file mode 100644
@@ -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 "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop auto worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop auto vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop auto gang vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop auto gang worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop auto worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop auto gang worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop auto worker
+ for (y = 0; y < 10; y++)
+#pragma acc loop auto vector
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop auto
+ for (y = 0; y < 10; y++)
+ ;
+
+#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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 loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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 loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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 loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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 loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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;
+}
new file mode 100644
@@ -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 "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop independent gang vector /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop independent gang worker /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop independent worker vector /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop independent gang worker vector /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (y = 0; y < 10; y++)
+ ;
+
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC worker loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+ }
+
+ return 0;
+}
new file mode 100644
@@ -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 "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ {
+ for (x = 0; x < 10; x++) /* { dg-message "note: 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;
+}
new file mode 100644
@@ -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 /* { dg-warning "region contains gang partitoned code but is not gang partitioned" } */
+ {
+ x = 0; /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } */
+ y = x < 10;
+ z = x++;
+ ;
+
+ y = 0;
+ z = y < 10;
+ x -= f_g (y++); /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */
+ ;
+
+ x = f_w (0); /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */
+ z = f_v (x < 10); /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ y -= f_s (x++); /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ ;
+
+ x = 0;
+ y = x < 10;
+ z = (x++);
+ y = 0;
+ x = y < 10;
+ z += (y++);
+ ;
+
+ x = 0;
+ y += f_s (x < 10); /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ x++;
+ y = 0;
+ y += f_v (y < 10); /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ y++;
+ z = 0;
+ y += f_w (z < 10); /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */
+ z++;
+ ;
+
+ x = 0;
+ y *= f_g ( /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */
+ f_w (x < 10) /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */
+ + f_g (x < 10) /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */
+ );
+ x++;
+ y = 0;
+ y *= y < 10;
+ y++;
+ z = 0;
+ y *= z < 10;
+ z++;
+ ;
+ }
+
+ return 0;
+}
new file mode 100644
@@ -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 "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop auto worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop auto vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop auto gang vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop auto gang worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop auto worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop auto gang worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop auto worker
+ for (y = 0; y < 10; y++)
+#pragma acc loop auto vector
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop auto
+ for (y = 0; y < 10; y++)
+ ;
+
+#pragma acc kernels loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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 loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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 loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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 loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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 loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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;
+}
new file mode 100644
@@ -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 "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop independent gang vector /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop independent gang worker /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop independent worker vector /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */
+ /* { dg-message "note: 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 "note: assigned OpenACC gang worker vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (y = 0; y < 10; y++)
+ ;
+
+#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC worker loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+ return 0;
+}
new file mode 100644
@@ -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 "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop seq
+ /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+ }
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent gang
+ /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+ }
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent worker
+ /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+ }
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent vector
+ /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+ }
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent gang vector
+ /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+ }
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent gang worker
+ /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+ }
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent worker vector
+ /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+ }
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent gang worker vector
+ /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+ }
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent gang
+ /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent
+ /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+ }
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent
+ /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent
+ /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop seq
+ /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent
+ /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop independent
+ /* { dg-message "note: 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 "note: assigned OpenACC seq loop parallelism" } */
+ /* Strangely indented to keep this similar to other test cases. */
+ if (c) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ {
+#pragma acc loop seq
+ /* { dg-message "note: 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;
+}
new file mode 100644
@@ -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
+#pragma acc loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto gang vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto gang worker /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto gang worker vector /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto gang /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop auto worker
+ for (y = 0; y < 10; y++)
+#pragma acc loop auto vector
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop auto
+ for (y = 0; y < 10; y++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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
+#pragma acc loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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
+#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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
+#pragma acc loop auto /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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
+#pragma acc loop /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" "" { target *-*-* } .-1 } */
+ 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;
+}
new file mode 100644
@@ -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
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent gang vector /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent gang worker /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent worker vector /* { dg-message "note: assigned OpenACC worker vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent gang worker vector /* { dg-message "note: assigned OpenACC gang worker vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent gang /* { dg-message "note: assigned OpenACC gang loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent worker /* { dg-message "note: assigned OpenACC worker loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent vector /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (y = 0; y < 10; y++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC worker loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang worker loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC vector loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ /* { dg-message "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } .-1 } */
+ for (x = 0; x < 10; x++)
+#pragma acc loop independent /* { dg-message "note: assigned OpenACC gang vector loop parallelism" } */
+ for (y = 0; y < 10; y++)
+#pragma acc loop seq /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (z = 0; z < 10; z++)
+ ;
+
+ return 0;
+}
new file mode 100644
@@ -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 "note: assigned OpenACC seq loop parallelism" } */
+ for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ ;
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ ;
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ for (y = 0; y < 10; y++)
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ ;
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ for (y = 0; y < 10; y++)
+ ;
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ for (y = 0; y < 10; y++)
+ for (z = 0; z < 10; z++)
+ ;
+
+#pragma acc kernels /* { dg-message "note: assigned OpenACC seq loop parallelism" } */
+ for (x = 0; x < 10; x++) /* { dg-message "note: beginning .parloops. region in OpenACC .kernels. construct" } */
+ for (y = 0; y < 10; y++)
+ for (z = 0; z < 10; z++)
+ ;
+
+ return 0;
+}
@@ -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" } */
@@ -91,7 +91,7 @@ extern void nohost (void);
int main ()
{
-#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32)
+#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32) /* { dg-warning "region contains gang partitoned code but is not gang partitioned" } */
{
gang ();
worker ();
@@ -21,12 +21,12 @@ 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_gangs(i) /* { dg-warning "is used uninitialized in this function" "TODO" { xfail *-*-* } } */
;
- #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized in this function" } */
+ #pragma acc kernels num_workers(j) /* { dg-warning "is used uninitialized in this function" "TODO" { xfail *-*-* } } */
;
- #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitialized in this function" } */
+ #pragma acc kernels vector_length(k) /* { dg-warning "is used uninitialized in this function" "TODO" { xfail *-*-* } } */
;
}
@@ -175,13 +175,13 @@ end subroutine sr5b
! { dg-final { scan-tree-dump-times "oacc_parallel device_type\\(\\*\\) \\\[ async\\(10\\) wait\\(10\\) num_gangs\\(10\\) num_workers\\(10\\) vector_length\\(10\\) \\\] device_type\\(nvidia_ptx\\) \\\[ async\\(3\\) wait\\(3\\) num_gangs\\(300\\) num_workers\\(300\\) vector_length\\(128\\) \\\] async\\(1\\) wait\\(1\\) num_gangs\\(1\\) num_workers\\(1\\) vector_length\\(1\\)" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ async\\(-1\\) wait\\(-1\\) \\\]" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single async\\(-1\\) num_gangs\\(1\\) device_type\\(nvidia\\) \\\[ async\\(-1\\) wait\\(-1\\) \\\]" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(nvidia\\) \\\[ async\\(1\\) wait\\(1\\) \\\]" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(nvidia\\) \\\[ async\\(1\\) wait\\(1\\) \\\]" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(\\*\\) \\\[ async\\(0\\) wait\\(0\\) \\\] device_type\\(nvidia\\) \\\[ async\\(2\\) wait\\(2\\) \\\]" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(\\*\\) \\\[ async\\(0\\) wait\\(0\\) \\\] device_type\\(nvidia\\) \\\[ async\\(2\\) wait\\(2\\) \\\]" 1 "omplower" } }
-! { dg-final { scan-tree-dump-times "oacc_kernels device_type\\(\\*\\) \\\[ async\\(0\\) wait\\(0\\) \\\] device_type\\(nvidia_ptx\\) \\\[ async\\(1\\) wait\\(1\\) \\\] async\\(-1\\) wait\\(-1\\)" 1 "omplower" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_gang_single num_gangs\\(1\\) device_type\\(\\*\\) \\\[ async\\(0\\) wait\\(0\\) \\\] device_type\\(nvidia_ptx\\) \\\[ async\\(1\\) wait\\(1\\) \\\] async\\(-1\\) wait\\(-1\\)" 1 "omplower" } }
! { dg-final { scan-tree-dump-times "acc loop device_type\\(nvidia\\) \\\[ tile\\(1\\) gang \\\] private\\(i1\\) private\\(i1\\.1\\)" 1 "omplower" } }
@@ -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_single async\\(-1\\)" 1 "convert_oacc_kernels" } }
+! { dg-final { scan-tree-dump-times "oacc_parallel_kernels_parallelized async\\(-1\\)" 1 "convert_oacc_kernels" } }
! { dg-final { scan-tree-dump-times "__builtin_GOACC_wait" 1 "convert_oacc_kernels" } }
@@ -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".
@@ -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
@@ -1,5 +1,21 @@
2019-01-31 Thomas Schwinge <thomas@codesourcery.com>
+ * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c:
+ Update.
+ * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-1.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/avoid-offloading-3.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
+ Likewise.
+ * testsuite/libgomp.oacc-fortran/avoid-offloading-1.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/avoid-offloading-2.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/avoid-offloading-3.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/initialize_kernels_loops.f90:
+ Likewise.
+
* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: New
file.
@@ -34,6 +34,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;
void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *event_info, acc_api_info *api_info)
{
@@ -50,7 +51,7 @@ void cb_enqueue_launch_start (acc_prof_info *prof_info, acc_event_info *event_in
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);
@@ -142,8 +143,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
@@ -172,8 +175,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;
@@ -181,6 +186,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" } */
{
@@ -206,8 +212,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;
@@ -215,6 +223,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 runtime setting" } */
{
@@ -1,5 +1,7 @@
/* Test that the compiler decides to "avoid offloading". */
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
#include <openacc.h>
int main(void)
@@ -7,8 +9,20 @@ int main(void)
int x, y;
#pragma acc data copyout(x, y)
-#pragma acc kernels /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } } */
- *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host);
+#pragma acc kernels
+ *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host); /* { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } */
+
+ /* The following will trigger "avoid offloading". */
+#pragma acc kernels
+ {
+#pragma acc loop auto /* { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" } */
+ /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } 18 } */
+ /* { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 18 } */
+ for (int i = 0; i < x; ++i)
+ if (x == 0)
+ x = 1;
+ ;
+ }
if (x != 33)
__builtin_abort();
@@ -1,6 +1,8 @@
/* Test that a user can override the compiler's "avoid offloading"
decision. */
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
#include <openacc.h>
int main(void)
@@ -19,8 +21,19 @@ int main(void)
int x, y;
#pragma acc data copyout(x, y)
-#pragma acc kernels /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } } */
- *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host);
+#pragma acc kernels
+ *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host); /* { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } */
+
+ /* The following will trigger "avoid offloading". */
+#pragma acc kernels
+ {
+#pragma acc loop auto /* { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" } */
+ /* { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } 30 } */
+ /* { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 30 } */
+ for (int i = 0; i < x; ++i)
+ if (x == 0)
+ x = 1;
+ }
if (x != 33)
__builtin_abort();
@@ -1,6 +1,8 @@
/* Test that a user can override the compiler's "avoid offloading"
decision. */
+/* { dg-additional-options "-fopt-info-optimized-omp" } */
+
/* Override the compiler's "avoid offloading" decision.
{ dg-additional-options "-foffload-force" } */
@@ -12,7 +14,17 @@ int main(void)
#pragma acc data copyout(x, y)
#pragma acc kernels
- *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host);
+ *((volatile int *) &x) = 33, y = acc_on_device (acc_device_host); /* { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } */
+
+ /* The following would trigger "avoid offloading". */
+#pragma acc kernels
+ {
+#pragma acc loop auto /* { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" } */
+ /* { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 22 } */
+ for (int i = 0; i < x; ++i)
+ if (x == 0)
+ x = 1;
+ }
if (x != 33)
__builtin_abort();
@@ -1,4 +1,3 @@
-/* { dg-additional-options "-fopenacc-kernels=split" } */
/* { dg-additional-options "-fopt-info-optimized-omp" } */
#undef NDEBUG
@@ -15,7 +14,7 @@ int main()
int c = 234; /* { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" } */
#pragma acc loop independent gang /* { dg-warning "note: assigned OpenACC gang loop parallelism" } */
- /* { dg-warning "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } 17 } */
+ /* { dg-warning "note: parallelized loop nest in OpenACC .kernels. construct" "" { target *-*-* } 16 } */
for (int i = 0; i < N; ++i)
b[i] = c;
@@ -2,6 +2,7 @@
! { dg-do run }
! { dg-additional-options "-cpp" }
+! { dg-additional-options "-fopt-info-optimized-omp" }
! As __OPTIMIZE__ is defined for -O1 and higher, we don't have an (easy) way to
! distinguish -O1 (where we will offload) from -O2 (where we won't offload), so
! for -O1 testing, we expect to abort.
@@ -10,16 +11,29 @@
IMPLICIT NONE
INCLUDE "openacc_lib.h"
+ INTEGER :: I
INTEGER, VOLATILE :: X
LOGICAL :: Y
!$ACC DATA COPYOUT(X, Y)
-!$ACC KERNELS ! { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } }
- X = 33
+!$ACC KERNELS
+ X = 33 ! { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" }
Y = ACC_ON_DEVICE (ACC_DEVICE_HOST);
!$ACC END KERNELS
!$ACC END DATA
+ ! The following will trigger "avoid offloading".
+!$ACC KERNELS
+!$ACC LOOP AUTO ! { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" }
+! { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } 27 }
+! { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 27 }
+ DO I = 1, X
+ IF (X .EQ. 0) THEN
+ X = 1
+ END IF
+ END DO
+!$ACC END KERNELS
+
IF (X .NE. 33) CALL ABORT
#if defined ACC_DEVICE_TYPE_nvidia
# if !defined __OPTIMIZE__
@@ -2,11 +2,13 @@
! { dg-do run }
! { dg-additional-options "-cpp" }
+! { dg-additional-options "-fopt-info-optimized-omp" }
IMPLICIT NONE
INCLUDE "openacc_lib.h"
INTEGER :: D
+ INTEGER :: I
INTEGER, VOLATILE :: X
LOGICAL :: Y
@@ -21,12 +23,24 @@
CALL ACC_INIT (D)
!$ACC DATA COPYOUT(X, Y)
-!$ACC KERNELS ! { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } }
- X = 33
+!$ACC KERNELS
+ X = 33 ! { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" }
Y = ACC_ON_DEVICE (ACC_DEVICE_HOST)
!$ACC END KERNELS
!$ACC END DATA
+ ! The following will trigger "avoid offloading".
+!$ACC KERNELS
+!$ACC LOOP AUTO ! { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" }
+! { dg-warning "OpenACC kernels construct will be executed sequentially; will by default avoid offloading to prevent data copy penalty" "" { target { openacc_nvidia_accel_selected && opt_levels_2_plus } } 34 }
+! { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 34 }
+ DO I = 1, X
+ IF (X .EQ. 0) THEN
+ X = 1
+ END IF
+ END DO
+!$ACC END KERNELS
+
IF (X .NE. 33) CALL ABORT
#if defined ACC_DEVICE_TYPE_nvidia
IF (Y) CALL ABORT
@@ -2,6 +2,7 @@
! { dg-do run }
! { dg-additional-options "-cpp" }
+! { dg-additional-options "-fopt-info-optimized-omp" }
! Override the compiler's "avoid offloading" decision.
! { dg-additional-options "-foffload-force" }
@@ -9,16 +10,28 @@
INCLUDE "openacc_lib.h"
INTEGER :: D
+ INTEGER :: I
INTEGER, VOLATILE :: X
LOGICAL :: Y
!$ACC DATA COPYOUT(X, Y)
!$ACC KERNELS
- X = 33
+ X = 33 ! { dg-warning "note: beginning .gang-single. region in OpenACC .kernels. construct" }
Y = ACC_ON_DEVICE (ACC_DEVICE_HOST)
!$ACC END KERNELS
!$ACC END DATA
+ ! The following would trigger "avoid offloading".
+!$ACC KERNELS
+!$ACC LOOP AUTO ! { dg-warning "note: forwarded loop nest in OpenACC .kernels. construct to .parloops. for analysis" }
+! { dg-warning "note: assigned OpenACC seq loop parallelism" "" { target *-*-* } 26 }
+ DO I = 1, X
+ IF (X .EQ. 0) THEN
+ X = 1
+ END IF
+ END DO
+!$ACC END KERNELS
+
IF (X .NE. 33) CALL ABORT
#if defined ACC_DEVICE_TYPE_nvidia
IF (Y) CALL ABORT
@@ -1,5 +1,4 @@
! { dg-do run }
-! { dg-additional-options "-fopenacc-kernels=split" }
! { dg-additional-options "-fopt-info-optimized-omp" }
subroutine kernel(lo, hi, a, b, c)
--
2.17.1