diff mbox

[committed,gomp4,3/3] Handle sequential code in kernels region - Testcases

Message ID 561BEDD1.7040203@mentor.com
State New
Headers show

Commit Message

Tom de Vries Oct. 12, 2015, 5:28 p.m. UTC 
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
diff mbox

Patch

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