@@ -1,4 +1,19 @@
2014-12-17 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * c-c++-common/goacc/asyncwait-1.c: New file.
+
+2014-12-17 Thomas Schwinge <thomas@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.
new file mode 100644
@@ -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;
+}
@@ -1,4 +1,4 @@
-/* { dg-do compile } */
+/* { dg-skip-if "not yet" { c++ } } */
int i, j, k;
extern int foo (void);
@@ -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 } */
;
}
@@ -1,3 +1,5 @@
+/* { dg-skip-if "not yet" { c++ } } */
+
void
f (void)
{
@@ -1,4 +1,4 @@
-/* { dg-do compile } */
+/* { dg-skip-if "not yet" { c++ } } */
int test1()
{
@@ -1,3 +1,5 @@
+/* { dg-skip-if "not yet" { c++ } } */
+
// pragma_external
#pragma acc update /* { dg-error "expected declaration specifiers before '#pragma'" } */
@@ -1,4 +1,4 @@
-// { dg-do compile }
+// { dg-skip-if "not yet" { c++ } }
void foo()
{
@@ -1,4 +1,4 @@
-// { dg-do compile }
+// { dg-skip-if "not yet" { c++ } }
void foo(int i)
{
@@ -1,4 +1,4 @@
-// { dg-do compile }
+// { dg-skip-if "not yet" { c++ } }
void f (void)
{
@@ -1,5 +1,41 @@
2014-12-17 Thomas Schwinge <thomas@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 <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+ Tom de Vries <tom@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 <jnorris@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 <thomas@codesourcery.com>
+
* testsuite/libgomp.oacc-c++/c++.exp: Don't add -flto to
ALWAYS_CFLAGS.
* testsuite/libgomp.oacc-c/c.exp: Likewise.
@@ -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
new file mode 100644
@@ -0,0 +1,17 @@
+/* { dg-do run } */
+/* { dg-shouldfail "" { *-*-* } { "*" } { "" } } */
+
+#include <stdlib.h>
+
+int
+main (void)
+{
+
+#pragma acc kernels
+ {
+ abort ();
+ }
+
+ return 0;
+}
+
new file mode 100644
@@ -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;
+}
+
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
@@ -1,4 +1,4 @@
-/* See also ../libgomp.oacc-c/collapse-4.c. */
+/* { dg-do run } */
#include <string.h>
@@ -1,5 +1,4 @@
/* { dg-do run } */
-/* { dg-additional-options "-std=c99" { target oacc_c } } */
#include <stdlib.h>
@@ -1,5 +1,4 @@
/* { dg-do run } */
-/* { dg-additional-options "-std=c99" { target oacc_c } } */
#include <stdlib.h>
new file mode 100644
@@ -0,0 +1,6 @@
+int
+main (void)
+{
+#pragma acc kernels
+ ;
+}
new file mode 100644
@@ -0,0 +1,6 @@
+int
+main (void)
+{
+#pragma acc parallel
+ ;
+}
deleted file mode 100644
@@ -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];
-}
new file mode 100644
@@ -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;
+}
@@ -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
deleted file mode 100644
@@ -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;
-}
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
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(-) Grüße, Thomas