Hi! On Thu, 16 Jul 2015 12:23:52 -0400, Nathan Sidwell 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 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 + Cesar Philippidis + + * testsuite/libgomp.oacc-c-c++-common/worker-partn-8.c: New file. + +2015-07-17 Thomas Schwinge * 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