Hi! Committed to gomp-4_0-branch in r218844:
commit 953f2d61152d9066365c8abb2e4551a17c4d0b83 Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4> Date: Wed Dec 17 22:36:40 2014 +0000 OpenACC testsuite updates. libgomp/ * testsuite/libgomp.oacc-c++/c++.exp (check_effective_target_oacc_c): Remove, and ... * testsuite/libgomp.oacc-c/c.exp (check_effective_target_oacc_c): ... likewise. Update all users. gcc/testsuite/ * c-c++-common/goacc/asyncwait-1.c: New file. libgomp/ * testsuite/libgomp.oacc-c-c++-common/abort-3.c: New file. * testsuite/libgomp.oacc-c-c++-common/abort-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/collapse-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/collapse-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-empty.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-empty.c: Likewise. * testsuite/libgomp.oacc-fortran/collapse-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/collapse-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/collapse-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/collapse-4.f90: Likewise. * testsuite/libgomp.oacc-fortran/collapse-5.f90: Likewise. * testsuite/libgomp.oacc-fortran/collapse-6.f90: Likewise. * testsuite/libgomp.oacc-fortran/collapse-7.f90: Likewise. * testsuite/libgomp.oacc-fortran/collapse-8.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/collapse-4.c: Update. * testsuite/libgomp.oacc-c/collapse-4.c: Remove. libgomp/ * testsuite/libgomp.oacc-c-c++-common/subr.h: New file. * testsuite/libgomp.oacc-c-c++-common/subr.cu: Remove. gcc/testsuite/ * c-c++-common/goacc/asyncwait-1.c: Disable for C++. * c-c++-common/goacc/collapse-1.c: Likewise. * c-c++-common/goacc/deviceptr-1.c: Likewise. * c-c++-common/goacc/if-clause-1.c: Likewise. * c-c++-common/goacc/loop-1.c: Likewise. * c-c++-common/goacc/pragma_context.c: Likewise. * c-c++-common/goacc/sb-1.c: Likewise. * c-c++-common/goacc/sb-2.c: Likewise. * c-c++-common/goacc/sb-3.c: Likewise. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@218844 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/testsuite/ChangeLog.gomp | 15 + gcc/testsuite/c-c++-common/goacc/asyncwait-1.c | 280 +++++++++++++ gcc/testsuite/c-c++-common/goacc/collapse-1.c | 2 +- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c | 22 +- gcc/testsuite/c-c++-common/goacc/if-clause-1.c | 2 + gcc/testsuite/c-c++-common/goacc/loop-1.c | 2 +- gcc/testsuite/c-c++-common/goacc/pragma_context.c | 2 + gcc/testsuite/c-c++-common/goacc/sb-1.c | 2 +- gcc/testsuite/c-c++-common/goacc/sb-2.c | 2 +- gcc/testsuite/c-c++-common/goacc/sb-3.c | 2 +- libgomp/ChangeLog.gomp | 36 ++ libgomp/testsuite/libgomp.oacc-c++/c++.exp | 4 - .../testsuite/libgomp.oacc-c-c++-common/abort-3.c | 17 + .../testsuite/libgomp.oacc-c-c++-common/abort-4.c | 17 + .../libgomp.oacc-c-c++-common/asyncwait-1.c | 466 +++++++++++++++++++++ .../libgomp.oacc-c-c++-common/collapse-1.c | 31 ++ .../libgomp.oacc-c-c++-common/collapse-2.c | 37 ++ .../libgomp.oacc-c-c++-common/collapse-3.c | 40 ++ .../libgomp.oacc-c-c++-common/collapse-4.c | 2 +- .../testsuite/libgomp.oacc-c-c++-common/data-2.c | 1 - .../testsuite/libgomp.oacc-c-c++-common/data-3.c | 1 - .../libgomp.oacc-c-c++-common/kernels-empty.c | 6 + .../libgomp.oacc-c-c++-common/parallel-empty.c | 6 + .../testsuite/libgomp.oacc-c-c++-common/subr.cu | 64 --- libgomp/testsuite/libgomp.oacc-c-c++-common/subr.h | 46 ++ libgomp/testsuite/libgomp.oacc-c/c.exp | 4 - libgomp/testsuite/libgomp.oacc-c/collapse-4.c | 27 -- .../testsuite/libgomp.oacc-fortran/collapse-1.f90 | 27 ++ .../testsuite/libgomp.oacc-fortran/collapse-2.f90 | 25 ++ .../testsuite/libgomp.oacc-fortran/collapse-3.f90 | 28 ++ .../testsuite/libgomp.oacc-fortran/collapse-4.f90 | 40 ++ .../testsuite/libgomp.oacc-fortran/collapse-5.f90 | 48 +++ .../testsuite/libgomp.oacc-fortran/collapse-6.f90 | 50 +++ .../testsuite/libgomp.oacc-fortran/collapse-7.f90 | 40 ++ .../testsuite/libgomp.oacc-fortran/collapse-8.f90 | 47 +++ 35 files changed, 1324 insertions(+), 117 deletions(-) diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index e8dcb44..dbedf73 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,4 +1,19 @@ 2014-12-17 Thomas Schwinge <tho...@codesourcery.com> + James Norris <jnor...@codesourcery.com> + + * c-c++-common/goacc/asyncwait-1.c: New file. + +2014-12-17 Thomas Schwinge <tho...@codesourcery.com> + + * c-c++-common/goacc/asyncwait-1.c: Disable for C++. + * c-c++-common/goacc/collapse-1.c: Likewise. + * c-c++-common/goacc/deviceptr-1.c: Likewise. + * c-c++-common/goacc/if-clause-1.c: Likewise. + * c-c++-common/goacc/loop-1.c: Likewise. + * c-c++-common/goacc/pragma_context.c: Likewise. + * c-c++-common/goacc/sb-1.c: Likewise. + * c-c++-common/goacc/sb-2.c: Likewise. + * c-c++-common/goacc/sb-3.c: Likewise. * c-c++-common/goacc-gomp/nesting-fail-1.c: Revise variable usage in constructs. diff --git gcc/testsuite/c-c++-common/goacc/asyncwait-1.c gcc/testsuite/c-c++-common/goacc/asyncwait-1.c new file mode 100644 index 0000000..9d69dcf --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/asyncwait-1.c @@ -0,0 +1,280 @@ +/* { dg-skip-if "not yet" { c++ } } */ + +#include <stdlib.h> + +int +main (int argc, char **argv) +{ + int N = 64; + float *a, *b; + int i; + + a = (float *) malloc (N * sizeof (float)); + b = (float *) malloc (N * sizeof (float)); + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1 2) /* { dg-error "expected '\\)' before numeric constant" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,) /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (,1) /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2,) /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2 3) /* { dg-error "expected '\\)' before numeric constant" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1,2,,) /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1 /* { dg-error "expected '\\)' before end of line" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (*) /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (a) + /* { dg-error "expected integer expression before" "" { target c } 85 } */ + /* { dg-error "'async' expression must be integral" "" { target c++ } 85 } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async (1.0) + /* { dg-error "expected integer expression before" "" { target c } 95 } */ + /* { dg-error "'async' expression must be integral" "" { target c++ } 95 } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async () /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) async + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1 2) /* { dg-error "expected '\\)' before numeric constant" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,) /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (,1) /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2,) /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2 3) /* { dg-error "expected '\\)' before numeric constant" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,2,,) /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1 /* { dg-error "expected '\\\)' before end of line" } */ + { + /* { dg-error "expected integer expression list before" "" { target c++ } 169 } */ + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,*) /* { dg-error "expected (primary-|)expression before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1,a) /*{ dg-error "must be integral" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (a) /* { dg-error "must be integral" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait (1.0) /* { dg-error "must be integral" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait () /* { dg-error "expected (integer |)expression (list |)before" } */ + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc parallel copyin (a[0:N]) copy (b[0:N]) wait + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc wait (1 2) /* { dg-error "expected '\\)' before numeric constant" } */ + +#pragma acc wait (1,) /* { dg-error "expected (primary-|)expression before" } */ + +#pragma acc wait (,1) /* { dg-error "expected (primary-|)expression before" } */ + +#pragma acc wait (1,2,) /* { dg-error "expected (primary-|)expression before" } */ + +#pragma acc wait (1,2 3) /* { dg-error "expected '\\)' before numeric constant" } */ + +#pragma acc wait (1,2,,) /* { dg-error "expected (primary-|)expression before" } */ + +#pragma acc wait (1 /* { dg-error "expected '\\\)' before end of line" } */ + /* { dg-error "expected integer expression list before" "" { target c++ } 238 } */ + +#pragma acc wait (1,*) /* { dg-error "expected (primary-|)expression before" } */ + +#pragma acc wait (1,a) /* { dg-error "expression must be integral" } */ + +#pragma acc wait (a) /* { dg-error "expression must be integral" } */ + +#pragma acc wait (1.0) /* { dg-error "expression must be integral" } */ + +#pragma acc wait 1 /* { dg-error "expected clause before numeric constant" } */ + +#pragma acc wait N /* { dg-error "expected clause before 'N'" } */ + +#pragma acc wait async (1 2) /* { dg-error "expected '\\)' before numeric constant" } */ + +#pragma acc wait async (1 2) /* { dg-error "expected '\\)' before numeric constant" } */ + +#pragma acc wait async (1,) /* { dg-error "expected (primary-|)expression before" } */ + +#pragma acc wait async (,1) /* { dg-error "expected (primary-|)expression before" } */ + +#pragma acc wait async (1,2,) /* { dg-error "expected (primary-|)expression before" } */ + +#pragma acc wait async (1,2 3) /* { dg-error "expected '\\)' before numeric constant" } */ + +#pragma acc wait async (1,2,,) /* { dg-error "expected (primary-|)expression before" } */ + +#pragma acc wait async (1 /* { dg-error "expected '\\)' before end of line" } */ + +#pragma acc wait async (*) /* { dg-error "expected (primary-|)expression before " } */ + +#pragma acc wait async (a) + /* { dg-error "expected integer expression before" "" { target c } 271 } */ + /* { dg-error "expression must be integral" "" { target c++ } 271 } */ + +#pragma acc wait async (1.0) + /* { dg-error "expected integer expression before" "" { target c } 275 } */ + /* { dg-error "expression must be integral" "" { target c++ } 275 } */ + + return 0; +} diff --git gcc/testsuite/c-c++-common/goacc/collapse-1.c gcc/testsuite/c-c++-common/goacc/collapse-1.c index b2aebb7..11b1438 100644 --- gcc/testsuite/c-c++-common/goacc/collapse-1.c +++ gcc/testsuite/c-c++-common/goacc/collapse-1.c @@ -1,4 +1,4 @@ -/* { dg-do compile } */ +/* { dg-skip-if "not yet" { c++ } } */ int i, j, k; extern int foo (void); diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-1.c gcc/testsuite/c-c++-common/goacc/deviceptr-1.c index cf2d809..546fa82 100644 --- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c +++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c @@ -1,3 +1,5 @@ +/* { dg-skip-if "not yet" { c++ } } */ + void fun1 (void) { @@ -9,24 +11,24 @@ fun1 (void) #pragma acc data deviceptr(fun1) /* { dg-error "'fun1' is not a variable" } */ ; #pragma acc parallel deviceptr(fun1[2:5]) - /* { dg-error "'fun1' is not a variable" "not a variable" { target *-*-* } 11 } */ - /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 11 } */ + /* { dg-error "'fun1' is not a variable" "not a variable" { target *-*-* } 13 } */ + /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 13 } */ ; int i; #pragma acc kernels deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */ ; #pragma acc data deviceptr(i[0:4]) - /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 19 } */ - /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 19 } */ + /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 21 } */ + /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 21 } */ ; float fa[10]; #pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */ ; #pragma acc kernels deviceptr(fa[1:5]) - /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 27 } */ - /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 27 } */ + /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 29 } */ + /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 29 } */ ; float *fp; @@ -42,10 +44,10 @@ fun2 (void) int i; float *fp; #pragma acc kernels deviceptr(fp,u,fun2,i,fp) - /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 44 } */ - /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 44 } */ - /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 44 } */ - /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 44 } */ + /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 46 } */ + /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 46 } */ + /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 46 } */ + /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 46 } */ ; } diff --git gcc/testsuite/c-c++-common/goacc/if-clause-1.c gcc/testsuite/c-c++-common/goacc/if-clause-1.c index c078596..85abf165 100644 --- gcc/testsuite/c-c++-common/goacc/if-clause-1.c +++ gcc/testsuite/c-c++-common/goacc/if-clause-1.c @@ -1,3 +1,5 @@ +/* { dg-skip-if "not yet" { c++ } } */ + void f (void) { diff --git gcc/testsuite/c-c++-common/goacc/loop-1.c gcc/testsuite/c-c++-common/goacc/loop-1.c index b890f38..fea40e0 100644 --- gcc/testsuite/c-c++-common/goacc/loop-1.c +++ gcc/testsuite/c-c++-common/goacc/loop-1.c @@ -1,4 +1,4 @@ -/* { dg-do compile } */ +/* { dg-skip-if "not yet" { c++ } } */ int test1() { diff --git gcc/testsuite/c-c++-common/goacc/pragma_context.c gcc/testsuite/c-c++-common/goacc/pragma_context.c index ad33d92..680dc9b 100644 --- gcc/testsuite/c-c++-common/goacc/pragma_context.c +++ gcc/testsuite/c-c++-common/goacc/pragma_context.c @@ -1,3 +1,5 @@ +/* { dg-skip-if "not yet" { c++ } } */ + // pragma_external #pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */ diff --git gcc/testsuite/c-c++-common/goacc/sb-1.c gcc/testsuite/c-c++-common/goacc/sb-1.c index bcb7272..5e55c95 100644 --- gcc/testsuite/c-c++-common/goacc/sb-1.c +++ gcc/testsuite/c-c++-common/goacc/sb-1.c @@ -1,4 +1,4 @@ -// { dg-do compile } +// { dg-skip-if "not yet" { c++ } } void foo() { diff --git gcc/testsuite/c-c++-common/goacc/sb-2.c gcc/testsuite/c-c++-common/goacc/sb-2.c index ec3eb95..a6760ec 100644 --- gcc/testsuite/c-c++-common/goacc/sb-2.c +++ gcc/testsuite/c-c++-common/goacc/sb-2.c @@ -1,4 +1,4 @@ -// { dg-do compile } +// { dg-skip-if "not yet" { c++ } } void foo(int i) { diff --git gcc/testsuite/c-c++-common/goacc/sb-3.c gcc/testsuite/c-c++-common/goacc/sb-3.c index 6c2926c..147b7b0 100644 --- gcc/testsuite/c-c++-common/goacc/sb-3.c +++ gcc/testsuite/c-c++-common/goacc/sb-3.c @@ -1,4 +1,4 @@ -// { dg-do compile } +// { dg-skip-if "not yet" { c++ } } void f (void) { diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index e0396b6..1fca220 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,5 +1,41 @@ 2014-12-17 Thomas Schwinge <tho...@codesourcery.com> + * testsuite/libgomp.oacc-c++/c++.exp + (check_effective_target_oacc_c): Remove, and ... + * testsuite/libgomp.oacc-c/c.exp (check_effective_target_oacc_c): + ... likewise. Update all users. + +2014-12-17 Thomas Schwinge <tho...@codesourcery.com> + James Norris <jnor...@codesourcery.com> + Cesar Philippidis <ce...@codesourcery.com> + Tom de Vries <t...@codesourcery.com> + + * testsuite/libgomp.oacc-c-c++-common/abort-3.c: New file. + * testsuite/libgomp.oacc-c-c++-common/abort-4.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/collapse-1.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/collapse-3.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/kernels-empty.c: Likewise. + * testsuite/libgomp.oacc-c-c++-common/parallel-empty.c: Likewise. + * testsuite/libgomp.oacc-fortran/collapse-1.f90: Likewise. + * testsuite/libgomp.oacc-fortran/collapse-2.f90: Likewise. + * testsuite/libgomp.oacc-fortran/collapse-3.f90: Likewise. + * testsuite/libgomp.oacc-fortran/collapse-4.f90: Likewise. + * testsuite/libgomp.oacc-fortran/collapse-5.f90: Likewise. + * testsuite/libgomp.oacc-fortran/collapse-6.f90: Likewise. + * testsuite/libgomp.oacc-fortran/collapse-7.f90: Likewise. + * testsuite/libgomp.oacc-fortran/collapse-8.f90: Likewise. + * testsuite/libgomp.oacc-c-c++-common/collapse-4.c: Update. + * testsuite/libgomp.oacc-c/collapse-4.c: Remove. + +2014-12-17 James Norris <jnor...@codesourcery.com> + + * testsuite/libgomp.oacc-c-c++-common/subr.h: New file. + * testsuite/libgomp.oacc-c-c++-common/subr.cu: Remove. + +2014-12-17 Thomas Schwinge <tho...@codesourcery.com> + * testsuite/libgomp.oacc-c++/c++.exp: Don't add -flto to ALWAYS_CFLAGS. * testsuite/libgomp.oacc-c/c.exp: Likewise. diff --git libgomp/testsuite/libgomp.oacc-c++/c++.exp libgomp/testsuite/libgomp.oacc-c++/c++.exp index f4d1f1a..2e2b618 100644 --- libgomp/testsuite/libgomp.oacc-c++/c++.exp +++ libgomp/testsuite/libgomp.oacc-c++/c++.exp @@ -13,10 +13,6 @@ if [info exists lang_include_flags] then { unset lang_include_flags } -proc check_effective_target_oacc_c { } { - return 0 -} - # Initialize dg. dg-init diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/abort-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/abort-3.c new file mode 100644 index 0000000..be7aaa8 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/abort-3.c @@ -0,0 +1,17 @@ +/* { dg-do run } */ +/* { dg-shouldfail "" { *-*-* } { "*" } { "" } } */ + +#include <stdlib.h> + +int +main (void) +{ + +#pragma acc kernels + { + abort (); + } + + return 0; +} + diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/abort-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/abort-4.c new file mode 100644 index 0000000..c29ca3f --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/abort-4.c @@ -0,0 +1,17 @@ +/* { dg-do run } */ + +#include <stdlib.h> + +int +main (int argc, char **argv) +{ + +#pragma acc kernels + { + if (argc != 1) + abort (); + } + + return 0; +} + diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c new file mode 100644 index 0000000..22cef6d --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c @@ -0,0 +1,466 @@ +/* { dg-do run { target openacc_nvidia_accel_selected } } */ +/* { dg-additional-options "-lcuda" } */ + +#include <openacc.h> +#include <stdlib.h> +#include "cuda.h" + +#include <stdio.h> +#include <sys/time.h> + +int +main (int argc, char **argv) +{ + CUresult r; + CUstream stream1; + int N = 128; //1024 * 1024; + float *a, *b, *c, *d, *e; + int i; + int nbytes; + + acc_init (acc_device_nvidia); + + nbytes = N * sizeof (float); + + a = (float *) malloc (nbytes); + b = (float *) malloc (nbytes); + c = (float *) malloc (nbytes); + d = (float *) malloc (nbytes); + e = (float *) malloc (nbytes); + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + } + +#pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N) + { + +#pragma acc parallel async + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc wait + + } + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 3.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 2.0; + b[i] = 0.0; + } + +#pragma acc data copy (a[0:N]) copy (b[0:N]) copyin (N) + { + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc wait (1) + + } + + for (i = 0; i < N; i++) + { + if (a[i] != 2.0) + abort (); + + if (b[i] != 2.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + } + +#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N) + { + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + } + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + } + + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; + } + +#pragma acc wait (1) + + } + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 9.0) + abort (); + + if (c[i] != 4.0) + abort (); + + if (d[i] != 1.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 2.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + e[i] = 0.0; + } + +#pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N) + { + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + } + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + } + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; + } + +#pragma acc parallel wait (1) async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; + } + +#pragma acc wait (1) + + } + + for (i = 0; i < N; i++) + { + if (a[i] != 2.0) + abort (); + + if (b[i] != 4.0) + abort (); + + if (c[i] != 4.0) + abort (); + + if (d[i] != 1.0) + abort (); + + if (e[i] != 11.0) + abort (); + } + + + r = cuStreamCreate (&stream1, CU_STREAM_NON_BLOCKING); + if (r != CUDA_SUCCESS) + { + fprintf (stderr, "cuStreamCreate failed: %d\n", r); + abort (); + } + + acc_set_cuda_stream (1, stream1); + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + b[i] = 0.0; + } + +#pragma acc data copy (a[0:N], b[0:N]) copyin (N) + { + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = a[ii]; + } + +#pragma acc wait (1) + + } + + for (i = 0; i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 5.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 7.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + } + +#pragma acc data copy (a[0:N]) copy (b[0:N]) copy (c[0:N]) copy (d[0:N]) copyin (N) + { + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + } + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + } + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; + } + +#pragma acc wait (1) + + } + + for (i = 0; i < N; i++) + { + if (a[i] != 7.0) + abort (); + + if (b[i] != 49.0) + abort (); + + if (c[i] != 4.0) + abort (); + + if (d[i] != 1.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + e[i] = 0.0; + } + +#pragma acc data copy (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) copyin (N) + { + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + } + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + } + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii]; + } + +#pragma acc parallel wait (1) async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + e[ii] = a[ii] + b[ii] + c[ii] + d[ii]; + } + +#pragma acc wait (1) + + } + + for (i = 0; i < N; i++) + { + if (a[i] != 3.0) + abort (); + + if (b[i] != 9.0) + abort (); + + if (c[i] != 4.0) + abort (); + + if (d[i] != 1.0) + abort (); + + if (e[i] != 17.0) + abort (); + } + + for (i = 0; i < N; i++) + { + a[i] = 4.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + e[i] = 0.0; + } + +#pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N) + { + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + } + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + } + +#pragma acc update host (a[0:N], b[0:N], c[0:N]) wait (1) + + } + + for (i = 0; i < N; i++) + { + if (a[i] != 4.0) + abort (); + + if (b[i] != 16.0) + abort (); + + if (c[i] != 4.0) + abort (); + } + + + for (i = 0; i < N; i++) + { + a[i] = 5.0; + b[i] = 0.0; + c[i] = 0.0; + d[i] = 0.0; + e[i] = 0.0; + } + +#pragma acc data copyin (a[0:N], b[0:N], c[0:N]) copyin (N) + { + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii]; + } + +#pragma acc parallel async (1) + { + int ii; + + for (ii = 0; ii < N; ii++) + c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii]; + } + +#pragma acc update host (a[0:N], b[0:N], c[0:N]) async (1) + +#pragma acc wait (1) + + } + + for (i = 0; i < N; i++) + { + if (a[i] != 5.0) + abort (); + + if (b[i] != 25.0) + abort (); + + if (c[i] != 4.0) + abort (); + } + + acc_shutdown (acc_device_nvidia); + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-1.c new file mode 100644 index 0000000..80fed6c --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-1.c @@ -0,0 +1,31 @@ +/* { dg-do run } */ + +#include <string.h> +#include <stdlib.h> + +int +main (void) +{ + int i, j, k, l = 0; + int a[3][3][3]; + + memset (a, '\0', sizeof (a)); + #pragma acc parallel + #pragma acc loop collapse(4 - 1) + for (i = 0; i < 2; i++) + for (j = 0; j < 2; j++) + for (k = 0; k < 2; k++) + a[i][j][k] = i + j * 4 + k * 16; + #pragma acc parallel + { + #pragma acc loop collapse(2) reduction(|:l) + for (i = 0; i < 2; i++) + for (j = 0; j < 2; j++) + for (k = 0; k < 2; k++) + if (a[i][j][k] != i + j * 4 + k * 16) + l = 1; + } + if (l) + abort (); + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-2.c new file mode 100644 index 0000000..44a77f7 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-2.c @@ -0,0 +1,37 @@ +/* { dg-do run } */ + +#include <stdlib.h> + +int +main (void) +{ + int i, j, k, l = 0, f = 0, x = 0; + int m1 = 4, m2 = -5, m3 = 17; + + #pragma acc parallel + #pragma acc loop collapse(3) reduction(+:l) + for (i = -2; i < m1; i++) + for (j = m2; j < -2; j++) + { + for (k = 13; k < m3; k++) + { + if ((i + 2) * 12 + (j + 5) * 4 + (k - 13) != 9 + f++) + l++; + } + } + + for (i = -2; i < m1; i++) + for (j = m2; j < -2; j++) + { + for (k = 13; k < m3; k++) + { + if ((i + 2) * 12 + (j + 5) * 4 + (k - 13) != 9 + f++) + x++; + } + } + + if (l != x) + abort (); + + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-3.c new file mode 100644 index 0000000..a5be728 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-3.c @@ -0,0 +1,40 @@ +/* { dg-do run } */ +/* { dg-options "-O2" } */ + +#include <string.h> +#include <stdlib.h> +#include <stdio.h> + +int +main (void) +{ + int i2, l = 0, r = 0; + int a[3][3][3]; + + memset (a, '\0', sizeof (a)); + #pragma acc parallel + #pragma acc loop collapse(4 - 1) + for (int i = 0; i < 2; i++) + for (int j = 0; j < 2; j++) + for (int k = 0; k < 2; k++) + a[i][j][k] = i + j * 4 + k * 16; +#pragma acc parallel + { + #pragma acc loop collapse(2) reduction(|:l) + for (i2 = 0; i2 < 2; i2++) + for (int j = 0; j < 2; j++) + for (int k = 0; k < 2; k++) + if (a[i2][j][k] != i2 + j * 4 + k * 16) + l += 1; + } + + for (i2 = 0; i2 < 2; i2++) + for (int j = 0; j < 2; j++) + for (int k = 0; k < 2; k++) + if (a[i2][j][k] != i2 + j * 4 + k * 16) + r += 1; + + if (l != r) + abort (); + return 0; +} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c index f3538a6..52dd435 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/collapse-4.c @@ -1,4 +1,4 @@ -/* See also ../libgomp.oacc-c/collapse-4.c. */ +/* { dg-do run } */ #include <string.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c index 3e02ffa..f867a66 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c @@ -1,5 +1,4 @@ /* { dg-do run } */ -/* { dg-additional-options "-std=c99" { target oacc_c } } */ #include <stdlib.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c index 44787dd..80c51d5 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-3.c @@ -1,5 +1,4 @@ /* { dg-do run } */ -/* { dg-additional-options "-std=c99" { target oacc_c } } */ #include <stdlib.h> diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-empty.c libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-empty.c new file mode 100644 index 0000000..a68a7cd --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-empty.c @@ -0,0 +1,6 @@ +int +main (void) +{ +#pragma acc kernels + ; +} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-empty.c libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-empty.c new file mode 100644 index 0000000..8e3bb43 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/parallel-empty.c @@ -0,0 +1,6 @@ +int +main (void) +{ +#pragma acc parallel + ; +} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/subr.cu libgomp/testsuite/libgomp.oacc-c-c++-common/subr.cu deleted file mode 100644 index e86e0fc..0000000 --- libgomp/testsuite/libgomp.oacc-c-c++-common/subr.cu +++ /dev/null @@ -1,64 +0,0 @@ - -extern "C" __global__ void -delay (clock_t * d_o, clock_t delay) -{ - clock_t start, ticks; - - start = clock (); - - ticks = 0; - - while (ticks < delay) - ticks = clock () - start; -} - -extern "C" __global__ void -delay2 (unsigned long *d_o, clock_t delay, unsigned long tid) -{ - clock_t start, ticks; - - start = clock (); - - ticks = 0; - - while (ticks < delay) - ticks = clock () - start; - - d_o[0] = tid; -} - -extern "C" __global__ void -sum (clock_t * d_o, int N) -{ - int i; - clock_t sum; - __shared__ clock_t ticks[32]; - - sum = 0; - - for (i = threadIdx.x; i < N; i += blockDim.x) - sum += d_o[i]; - - ticks[threadIdx.x] = sum; - - syncthreads (); - - for (i = 16; i >= 1; i >>= 1) - { - if (threadIdx.x < i) - ticks[threadIdx.x] += ticks[threadIdx.x + i]; - - syncthreads (); - } - - d_o[0] = ticks[0]; -} - -extern "C" __global__ void -mult (int n, float *x, float *y) -{ - int i = blockIdx.x * blockDim.x + threadIdx.x; - - for (i = 0; i < n; i++) - y[i] = x[i] * x[i]; -} diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/subr.h libgomp/testsuite/libgomp.oacc-c-c++-common/subr.h new file mode 100644 index 0000000..9db236c --- /dev/null +++ libgomp/testsuite/libgomp.oacc-c-c++-common/subr.h @@ -0,0 +1,46 @@ + +#if ACC_DEVICE_TYPE_nvidia + +#pragma acc routine nohost +static int clock (void) +{ + int thetime; + + asm __volatile__ ("mov.u32 %0, %%clock;" : "=r"(thetime)); + + return thetime; +} + +#endif + +void +delay (unsigned long *d_o, unsigned long delay) +{ + int start, ticks; + + start = clock (); + + ticks = 0; + + while (ticks < delay) + ticks = clock () - start; + + return; +} + +void +delay2 (unsigned long *d_o, unsigned long delay, unsigned long tid) +{ + int start, ticks; + + start = clock (); + + ticks = 0; + + while (ticks < delay) + ticks = clock () - start; + + d_o[0] = tid; + + return; +} diff --git libgomp/testsuite/libgomp.oacc-c/c.exp libgomp/testsuite/libgomp.oacc-c/c.exp index 2e973c5..d7e5536 100644 --- libgomp/testsuite/libgomp.oacc-c/c.exp +++ libgomp/testsuite/libgomp.oacc-c/c.exp @@ -19,10 +19,6 @@ if ![info exists DEFAULT_CFLAGS] then { set DEFAULT_CFLAGS "-O2" } -proc check_effective_target_oacc_c { } { - return 1 -} - # Initialize dg. dg-init diff --git libgomp/testsuite/libgomp.oacc-c/collapse-4.c libgomp/testsuite/libgomp.oacc-c/collapse-4.c deleted file mode 100644 index 17663ca..0000000 --- libgomp/testsuite/libgomp.oacc-c/collapse-4.c +++ /dev/null @@ -1,27 +0,0 @@ -/* See also ../libgomp.oacc-c-c++-common/collapse-4.c. */ -/* { dg-options "-O2 -std=c99" } */ - -#include <string.h> - -int -main (void) -{ - int l = 0; - int b[3][3]; - - memset (b, '\0', sizeof (b)); - -#pragma acc parallel copy(b[0:3][0:3]) copy(l) - { -#pragma acc loop collapse(2) reduction(+:l) - for (int i = 0; i < 2; i++) - for (int j = 0; j < 2; j++) - if (b[i][j] != 16) - l += 1; - } - - if (l != 2 * 2) - __builtin_abort(); - - return 0; -} diff --git libgomp/testsuite/libgomp.oacc-fortran/collapse-1.f90 libgomp/testsuite/libgomp.oacc-fortran/collapse-1.f90 new file mode 100644 index 0000000..4c07bc2 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/collapse-1.f90 @@ -0,0 +1,27 @@ +! { dg-do run } + +program collapse1 + integer :: i, j, k, a(1:3, 4:6, 5:7) + logical :: l + l = .false. + a(:, :, :) = 0 + !$acc parallel + !$acc loop collapse(4 - 1) + do i = 1, 3 + do j = 4, 6 + do k = 5, 7 + a(i, j, k) = i + j + k + end do + end do + end do + !$acc loop collapse(2) reduction(.or.:l) + do i = 1, 3 + do j = 4, 6 + do k = 5, 7 + if (a(i, j, k) .ne. (i + j + k)) l = .true. + end do + end do + end do + !$acc end parallel + if (l) call abort +end program collapse1 diff --git libgomp/testsuite/libgomp.oacc-fortran/collapse-2.f90 libgomp/testsuite/libgomp.oacc-fortran/collapse-2.f90 new file mode 100644 index 0000000..ca3b638 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/collapse-2.f90 @@ -0,0 +1,25 @@ +! { dg-do run } + +program collapse2 + integer :: i, j, k, a(1:3, 4:6, 5:7) + logical :: l + l = .false. + a(:, :, :) = 0 + !$acc parallel + !$acc loop collapse(4 - 1) + do 164 i = 1, 3 + do 164 j = 4, 6 + do 164 k = 5, 7 + a(i, j, k) = i + j + k +164 end do + !$acc loop collapse(2) reduction(.or.:l) +firstdo: do i = 1, 3 + do j = 4, 6 + do k = 5, 7 + if (a(i, j, k) .ne. (i + j + k)) l = .true. + end do + end do + end do firstdo + !$acc end parallel + if (l) call abort +end program collapse2 diff --git libgomp/testsuite/libgomp.oacc-fortran/collapse-3.f90 libgomp/testsuite/libgomp.oacc-fortran/collapse-3.f90 new file mode 100644 index 0000000..50e6100 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/collapse-3.f90 @@ -0,0 +1,28 @@ +! { dg-do run } + +program collapse3 + integer :: a(3,3,3), k, kk, kkk, l, ll, lll + !$acc parallel + !$acc loop collapse(3) + do 115 k=1,3 +dokk: do kk=1,3 + do kkk=1,3 + a(k,kk,kkk) = 1 + enddo + enddo dokk +115 continue + !$acc end parallel + if (any(a(1:3,1:3,1:3).ne.1)) call abort + + !$acc parallel + !$acc loop collapse(3) +dol: do 120 l=1,3 +doll: do ll=1,3 + do lll=1,3 + a(l,ll,lll) = 2 + enddo + enddo doll +120 end do dol + !$acc end parallel + if (any(a(1:3,1:3,1:3).ne.2)) call abort +end program collapse3 diff --git libgomp/testsuite/libgomp.oacc-fortran/collapse-4.f90 libgomp/testsuite/libgomp.oacc-fortran/collapse-4.f90 new file mode 100644 index 0000000..41b66db --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/collapse-4.f90 @@ -0,0 +1,40 @@ +! { dg-do run } + +! collapse3.f90:test1 +program collapse4 + integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19) + logical :: l, r + l = .false. + r = .false. + a(:, :, :) = 0 + b(:, :, :) = 0 + !$acc parallel + !$acc loop collapse (3) reduction (.or.:l) + do i = 2, 6 + do j = -2, 4 + do k = 13, 18 + l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4 + l = l.or.k.lt.13.or.k.gt.18 + if (.not.l) a(i, j, k) = a(i, j, k) + 1 + end do + end do + end do + !$acc end parallel + do i = 2, 6 + do j = -2, 4 + do k = 13, 18 + r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4 + r = r.or.k.lt.13.or.k.gt.18 + if (.not.l) b(i, j, k) = b(i, j, k) + 1 + end do + end do + end do + if (l .neqv. r) call abort + do i = 2, 6 + do j = -2, 4 + do k = 13, 18 + if (a(i, j, k) .ne. b(i, j, k)) call abort + end do + end do + end do +end program collapse4 diff --git libgomp/testsuite/libgomp.oacc-fortran/collapse-5.f90 libgomp/testsuite/libgomp.oacc-fortran/collapse-5.f90 new file mode 100644 index 0000000..8c20f04 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/collapse-5.f90 @@ -0,0 +1,48 @@ +! { dg-do run } + +! collapse3.f90:test2 +program collapse5 + integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19) + integer :: v1, v2, v3, v4, v5, v6 + logical :: l, r + l = .false. + r = .false. + a(:, :, :) = 0 + b(:, :, :) = 0 + v1 = 3 + v2 = 6 + v3 = -2 + v4 = 4 + v5 = 13 + v6 = 18 + !$acc parallel + !$acc loop collapse (3) reduction (.or.:l) + do i = v1, v2 + do j = v3, v4 + do k = v5, v6 + l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4 + l = l.or.k.lt.13.or.k.gt.18 + if (.not.l) a(i, j, k) = a(i, j, k) + 1 + m = i * 100 + j * 10 + k + end do + end do + end do + !$acc end parallel + do i = v1, v2 + do j = v3, v4 + do k = v5, v6 + r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4 + r = r.or.k.lt.13.or.k.gt.18 + if (.not.l) b(i, j, k) = b(i, j, k) + 1 + end do + end do + end do + if (l .neqv. r) call abort + do i = v1, v2 + do j = v3, v4 + do k = v5, v6 + if (a(i, j, k) .ne. b(i, j, k)) call abort + end do + end do + end do +end program collapse5 diff --git libgomp/testsuite/libgomp.oacc-fortran/collapse-6.f90 libgomp/testsuite/libgomp.oacc-fortran/collapse-6.f90 new file mode 100644 index 0000000..7404b91 --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/collapse-6.f90 @@ -0,0 +1,50 @@ +! { dg-do run } + +! collapse3.f90:test3 +program collapse6 + integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19) + integer :: v1, v2, v3, v4, v5, v6, v7, v8, v9 + logical :: l, r + l = .false. + r = .false. + a(:, :, :) = 0 + b(:, :, :) = 0 + v1 = 3 + v2 = 6 + v3 = -2 + v4 = 4 + v5 = 13 + v6 = 18 + v7 = 1 + v8 = 1 + v9 = 1 + !$acc parallel + !$acc loop collapse (3) reduction (.or.:l) + do i = v1, v2, v7 + do j = v3, v4, v8 + do k = v5, v6, v9 + l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4 + l = l.or.k.lt.13.or.k.gt.18 + if (.not.l) a(i, j, k) = a(i, j, k) + 1 + end do + end do + end do + !$acc end parallel + do i = v1, v2, v7 + do j = v3, v4, v8 + do k = v5, v6, v9 + r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4 + r = r.or.k.lt.13.or.k.gt.18 + if (.not.r) b(i, j, k) = b(i, j, k) + 1 + end do + end do + end do + if (l .neqv. r) call abort + do i = v1, v2, v7 + do j = v3, v4, v8 + do k = v5, v6, v9 + if (a(i, j, k) .ne. b(i, j, k)) call abort + end do + end do + end do +end program collapse6 diff --git libgomp/testsuite/libgomp.oacc-fortran/collapse-7.f90 libgomp/testsuite/libgomp.oacc-fortran/collapse-7.f90 new file mode 100644 index 0000000..12efd8c --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/collapse-7.f90 @@ -0,0 +1,40 @@ +! { dg-do run } + +! collapse3.f90:test4 +program collapse7 + integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19) + logical :: l, r + l = .false. + r = .false. + a(:, :, :) = 0 + b(:, :, :) = 0 + !$acc parallel + !$acc loop collapse (3) reduction (.or.:l) + do i = 2, 6 + do j = -2, 4 + do k = 13, 18 + l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4 + l = l.or.k.lt.13.or.k.gt.18 + if (.not.l) a(i, j, k) = a(i, j, k) + 1 + end do + end do + end do + !$acc end parallel + do i = 2, 6 + do j = -2, 4 + do k = 13, 18 + r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4 + r = r.or.k.lt.13.or.k.gt.18 + if (.not.r) b(i, j, k) = b(i, j, k) + 1 + end do + end do + end do + if (l .neqv. r) call abort + do i = 1, 7 + do j = -3, 5 + do k = 12, 19 + if (a(i, j, k) .ne. b(i, j, k)) call abort + end do + end do + end do +end program collapse7 diff --git libgomp/testsuite/libgomp.oacc-fortran/collapse-8.f90 libgomp/testsuite/libgomp.oacc-fortran/collapse-8.f90 new file mode 100644 index 0000000..04fbcfe --- /dev/null +++ libgomp/testsuite/libgomp.oacc-fortran/collapse-8.f90 @@ -0,0 +1,47 @@ +! { dg-do run } + +! collapse3.f90:test5 +program collapse8 + integer :: i, j, k, a(1:7, -3:5, 12:19), b(1:7, -3:5, 12:19) + integer :: v1, v2, v3, v4, v5, v6 + logical :: l, r + l = .false. + r = .false. + a(:, :, :) = 0 + b(:, :, :) = 0 + v1 = 3 + v2 = 6 + v3 = -2 + v4 = 4 + v5 = 13 + v6 = 18 + !$acc parallel + !$acc loop collapse (3) reduction (.or.:l) + do i = v1, v2 + do j = v3, v4 + do k = v5, v6 + l = l.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4 + l = l.or.k.lt.13.or.k.gt.18 + if (.not.l) a(i, j, k) = a(i, j, k) + 1 + end do + end do + end do + !$acc end parallel + do i = v1, v2 + do j = v3, v4 + do k = v5, v6 + r = r.or.i.lt.2.or.i.gt.6.or.j.lt.-2.or.j.gt.4 + r = r.or.k.lt.13.or.k.gt.18 + if (.not.r) b(i, j, k) = b(i, j, k) + 1 + end do + end do + end do + if (l .neqv. r) call abort + do i = v1, v2 + do j = v3, v4 + do k = v5, v6 + if (a(i, j, k) .ne. b(i, j, k)) call abort + end do + end do + end do +end program collapse8 Grüße, Thomas
pgpffgXJfBIVb.pgp
Description: PGP signature