aboutsummaryrefslogtreecommitdiff
path: root/libgomp
diff options
context:
space:
mode:
authorTom de Vries <tdevries@suse.de>2018-12-19 14:20:44 +0000
committerTom de Vries <vries@gcc.gnu.org>2018-12-19 14:20:44 +0000
commit49188cd1f2deb943e4047dbffba7a333875d6479 (patch)
treee2a4813813fd016575e3cbc5e7f561a7c933b117 /libgomp
parent9b09e453c6884f2ea391dc70c188c159f0883150 (diff)
downloadgcc-49188cd1f2deb943e4047dbffba7a333875d6479.zip
gcc-49188cd1f2deb943e4047dbffba7a333875d6479.tar.gz
gcc-49188cd1f2deb943e4047dbffba7a333875d6479.tar.bz2
[nvptx, libgomp] Move rtl-dump test-cases to libgomp
The goacc.exp test-cases nvptx-merged-loop.c and nvptx-sese-1.c are failing during linking due to missing libgomp.spec. Move them to the libgomp testsuite. Build and reg-tested on x86_64 with nvptx accelerator. 2018-12-19 Tom de Vries <tdevries@suse.de> * gcc.dg/goacc/nvptx-merged-loop.c: Move to libgomp/testsuite/libgomp.oacc-c-c++-common. * gcc.dg/goacc/nvptx-sese-1.c: Same. * testsuite/lib/libgomp.exp: Add load_lib of scanoffloadrtl.exp. * testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c: Move from gcc/testsuite/gcc.dg/goacc. * testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Same. From-SVN: r267267
Diffstat (limited to 'libgomp')
-rw-r--r--libgomp/ChangeLog7
-rw-r--r--libgomp/testsuite/lib/libgomp.exp1
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c30
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c35
4 files changed, 73 insertions, 0 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 084f174..ad0abb8 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,10 @@
+2018-12-19 Tom de Vries <tdevries@suse.de>
+
+ * testsuite/lib/libgomp.exp: Add load_lib of scanoffloadrtl.exp.
+ * testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c: Move from
+ gcc/testsuite/gcc.dg/goacc.
+ * testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Same.
+
2018-12-14 Thomas Schwinge <thomas@codesourcery.com>
Chung-Lin Tang <cltang@codesourcery.com>
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index c41b3e6..04738a9 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -31,6 +31,7 @@ load_gcc_lib scanrtl.exp
load_gcc_lib scantree.exp
load_gcc_lib scanltranstree.exp
load_gcc_lib scanoffloadtree.exp
+load_gcc_lib scanoffloadrtl.exp
load_gcc_lib scanipa.exp
load_gcc_lib scanwpaipa.exp
load_gcc_lib timeout-dg.exp
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c
new file mode 100644
index 0000000..8a2117e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c
@@ -0,0 +1,30 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-options "-foffload=-fdump-rtl-mach" } */
+/* { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
+
+#define N (32*32*32+17)
+void __attribute__ ((noinline)) Foo (int *ary)
+{
+ int ix;
+
+#pragma acc parallel num_workers(32) vector_length(32) copyout(ary[0:N])
+ {
+ /* Loop partitioning should be merged. */
+#pragma acc loop worker vector
+ for (unsigned ix = 0; ix < N; ix++)
+ {
+ ary[ix] = ix;
+ }
+ }
+}
+
+int main ()
+{
+ int ary[N];
+
+ Foo (ary);
+
+ return 0;
+}
+
+/* { dg-final { scan-offload-rtl-dump "Merging loop .* into " "mach" } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c
new file mode 100644
index 0000000..9583265
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c
@@ -0,0 +1,35 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-options "-foffload=-fdump-rtl-mach" } */
+/* { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
+
+#pragma acc routine seq
+int __attribute__((noinline)) foo (int x)
+{
+ return x & 2;
+}
+
+int main ()
+{
+ int r = 0;
+
+#pragma acc parallel copy(r) vector_length(32)
+ {
+#pragma acc loop vector reduction (+:r)
+ for (int i = 00; i < 40; i++)
+ r += i;
+
+ /* This piece is a multi-block SESE region */
+ if (foo (r))
+ r *= 2;
+
+ if (r & 1) /* to here. */
+#pragma acc loop vector reduction (+:r)
+ for (int i = 00; i < 40; i++)
+ r += i;
+ }
+
+ return 0;
+}
+
+/* Match {N->N(.N)+} */
+/* { dg-final { scan-offload-rtl-dump "SESE regions:.* \[0-9\]+{\[0-9\]+->\[0-9\]+(\\.\[0-9\]+)+}" "mach" } } */