Hi! On Thu, 16 Jul 2015 12:23:52 -0400, Nathan Sidwell <nat...@acm.org> wrote: > I've committed this patch to fix a bug in the worker spill/fill code. We > ended > up not incrementing the pointer, resulting in the stack frame being filled > with > the same value. > > Thanks to Jim for finding the failure.
Cesar had prepared a reduced test case, a slightly altered variant of
which I've now committed to gomp-4_0-branch in r225924:
commit ee7fb343a0d0dbd17ac8dc7d24048d8647e41232
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri Jul 17 09:11:10 2015 +0000
OpenACC: Add test case for worker state propagation handling the stack frame
... for problem that got addressed in r225896.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/worker-partn-8.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@225924
138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog.gomp | 5 ++
.../libgomp.oacc-c-c++-common/worker-partn-8.c | 53 ++++++++++++++++++++++
2 files changed, 58 insertions(+)
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 0293ad5..ec943f5 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,4 +1,9 @@
2015-07-17 Thomas Schwinge <tho...@codesourcery.com>
+ Cesar Philippidis <ce...@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/worker-partn-8.c: New file.
+
+2015-07-17 Thomas Schwinge <tho...@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-5.c:
Remove XFAIL.
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/worker-partn-8.c
libgomp/testsuite/libgomp.oacc-c-c++-common/worker-partn-8.c
new file mode 100644
index 0000000..e787947
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/worker-partn-8.c
@@ -0,0 +1,53 @@
+/* { dg-additional-options "-O0" } */
+
+/* With -O0, variables are on the stack, not in registers. Check that worker
+ state propagation handles the stack frame. */
+
+int
+main (int argc, char *argv[])
+{
+ int w0 = 0;
+ int w1 = 0;
+ int w2 = 0;
+ int w3 = 0;
+ int w4 = 0;
+ int w5 = 0;
+ int w6 = 0;
+ int w7 = 0;
+
+ int i;
+
+#pragma acc parallel num_gangs (1) num_workers (8) copy (w0, w1, w2, w3, w4,
w5, w6, w7)
+ {
+ int internal = 100;
+
+#pragma acc loop worker
+ for (i = 0; i < 8; i++)
+ {
+ switch (i)
+ {
+ case 0: w0 = internal; break;
+ case 1: w1 = internal; break;
+ case 2: w2 = internal; break;
+ case 3: w3 = internal; break;
+ case 4: w4 = internal; break;
+ case 5: w5 = internal; break;
+ case 6: w6 = internal; break;
+ case 7: w7 = internal; break;
+ default: break;
+ }
+ }
+ }
+
+ if (w0 != 100
+ || w1 != 100
+ || w2 != 100
+ || w3 != 100
+ || w4 != 100
+ || w5 != 100
+ || w6 != 100
+ || w7 != 100)
+ __builtin_abort ();
+
+ return 0;
+}
Grüße,
Thomas
signature.asc
Description: PGP signature
