This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[committed, gomp4, 3/3] Handle sequential code in kernels region - Testcases
- From: Tom de Vries <Tom_deVries at mentor dot com>
- To: "gcc-patches at gnu dot org" <gcc-patches at gnu dot org>
- Cc: Jakub Jelinek <jakub at redhat dot com>
- Date: Mon, 12 Oct 2015 19:28:49 +0200
- Subject: [committed, gomp4, 3/3] Handle sequential code in kernels region - Testcases
- Authentication-results: sourceware.org; auth=none
- References: <561BEA02 dot 6010808 at mentor dot com>
On 12/10/15 19:12, Tom de Vries wrote:
Hi,
I've committed the following patch series.
1 Add get_bbs_in_oacc_kernels_region
2 Handle sequential code in kernels region
3 Handle sequential code in kernels region - Testcases
The patch series adds detection of whether sequential code (that is,
code in the oacc kernels region before and after the loop that is to be
parallelized), is safe to execute in parallel.
Bootstrapped and reg-tested on x86_64.
I'll post the patches individually, in reply to this email.
This patch adds relevant test-cases.
Thanks,
- Tom
Handle sequential code in kernels region - Testcases
2015-10-12 Tom de Vries <tom@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c: New test.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c: New test.
* testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c: New test.
---
.../kernels-loop-and-seq-2.c | 36 +++++++++++++++++++++
.../kernels-loop-and-seq-3.c | 37 ++++++++++++++++++++++
.../kernels-loop-and-seq-4.c | 36 +++++++++++++++++++++
.../kernels-loop-and-seq-5.c | 37 ++++++++++++++++++++++
.../kernels-loop-and-seq-6.c | 36 +++++++++++++++++++++
.../kernels-loop-and-seq.c | 37 ++++++++++++++++++++++
6 files changed, 219 insertions(+)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c
new file mode 100644
index 0000000..2e4100f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-2.c
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+#pragma acc kernels copy (a[0:N])
+ {
+ a[0] = a[0] + 1;
+
+ for (int i = 0; i < n; i++)
+ a[i] = 1;
+ }
+
+ return a[0];
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 1)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c
new file mode 100644
index 0000000..b3e736b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-3.c
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+
+#pragma acc kernels copy (a[0:N])
+ {
+ for (int i = 0; i < n; i++)
+ a[i] = 1;
+
+ a[0] = 2;
+ }
+
+ return a[0];
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 2)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c
new file mode 100644
index 0000000..8b9affa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-4.c
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+#pragma acc kernels copy (a[0:N])
+ {
+ a[0] = 2;
+
+ for (int i = 0; i < n; i++)
+ a[i] = 1;
+ }
+
+ return a[0];
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 1)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c
new file mode 100644
index 0000000..83d4e7f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-5.c
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+ int r;
+#pragma acc kernels copyout(r) copy (a[0:N])
+ {
+ r = a[0];
+
+ for (int i = 0; i < n; i++)
+ a[i] = 1;
+ }
+
+ return r;
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 0)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c
new file mode 100644
index 0000000..01d5e5e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq-6.c
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+#pragma acc kernels copy (a[0:N])
+ {
+ int r = a[0];
+
+ for (int i = 0; i < n; i++)
+ a[i] = 1 + r;
+ }
+
+ return a[0];
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 1)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c
new file mode 100644
index 0000000..61d1283
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-and-seq.c
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+/* { dg-additional-options "-ftree-parallelize-loops=32" } */
+
+#include <stdlib.h>
+
+#define N 32
+
+unsigned int
+foo (int n, unsigned int *a)
+{
+
+#pragma acc kernels copy (a[0:N])
+ {
+ for (int i = 0; i < n; i++)
+ a[i] = 1;
+
+ a[0] = a[0] + 1;
+ }
+
+ return a[0];
+}
+
+int
+main (void)
+{
+ unsigned int a[N];
+ unsigned res, i;
+
+ for (i = 0; i < N; ++i)
+ a[i] = i % 4;
+
+ res = foo (N, a);
+ if (res != 2)
+ abort ();
+
+ return 0;
+}
--
1.9.1