* [gomp4] Tests for private variables/state propagation
@ 2015-06-17 14:23 Julian Brown
2015-06-18 14:33 ` Thomas Schwinge
0 siblings, 1 reply; 2+ messages in thread
From: Julian Brown @ 2015-06-17 14:23 UTC (permalink / raw)
To: gcc-patches, nathan, Cesar Philippidis, Thomas Schwinge, Bernd Schmidt
[-- Attachment #1: Type: text/plain, Size: 1068 bytes --]
Hi,
This is a set of tests for OpenACC private variable/state propagation
support in GCC. The associated functionality is a work-in-progress: as
such, many of these tests do not pass yet (causing incorrect results,
ICEs or even bogus assembly output). I believe the tests to be valid
OpenACC, though it's possible I misinterpreted the spec at some points!
I will apply to the gomp4 branch shortly. (We will of course be working
on addressing the failures.)
Cheers,
Julian
ChangeLog
libgomp/
* testsuite/libgomp.oacc-c-c++-common/
private-vars-par-gang-{1,2,3}.c: New tests.
* testsuite/libgomp.oacc-c-c++-common/
private-vars-local-gang-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/
private-vars-loop-gang-{1,2,3,4,5,6}.c: New tests.
* testsuite/libgomp.oacc-c-c++-common/
private-vars-loop-worker-{1,2,3,4,5,6,7}.c: New tests.
* testsuite/libgomp.oacc-c-c++-common/
private-vars-local-worker-{1,2,3,4,5}.c: New tests.
* testsuite/libgomp.oacc-c-c++-common/
private-vars-loop-vector-{1,2}.c: New tests.
[-- Attachment #2: private-var-tests-1.diff --]
[-- Type: text/x-patch, Size: 30870 bytes --]
commit 40193f49480f0a0b750d15049d29fd427282c5f0
Author: Julian Brown <julian@codesourcery.com>
Date: Tue Jun 16 03:50:55 2015 -0700
New set of private variable/state propagation tests.
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-gang-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-gang-1.c
new file mode 100644
index 0000000..ada46d0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-gang-1.c
@@ -0,0 +1,38 @@
+#include <assert.h>
+
+/* Test of gang-private variables declared in local scope with parallel
+ directive. */
+
+#if defined(ACC_DEVICE_TYPE_host) || defined(ACC_DEVICE_TYPE_host_nonshm)
+#define ACTUAL_GANGS 1
+#else
+#define ACTUAL_GANGS 32
+#endif
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[ACTUAL_GANGS];
+
+ for (i = 0; i < ACTUAL_GANGS; i++)
+ arr[i] = 3;
+
+ #pragma acc parallel copy(arr) num_gangs(ACTUAL_GANGS) num_workers(8) \
+ vector_length(32)
+ {
+ int x;
+
+ #pragma acc loop gang(static:1)
+ for (i = 0; i < ACTUAL_GANGS; i++)
+ x = i * 2;
+
+ #pragma acc loop gang(static:1)
+ for (i = 0; i < ACTUAL_GANGS; i++)
+ arr[i] += x;
+ }
+
+ for (i = 0; i < ACTUAL_GANGS; i++)
+ assert (arr[i] == 3 + i * 2);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-1.c
new file mode 100644
index 0000000..f8658e5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-1.c
@@ -0,0 +1,56 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared in a local scope, broadcasting
+ to vector-partitioned mode. Back-to-back worker loops. */
+
+int
+main (int argc, char* argv[])
+{
+ int i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ int x = i ^ j * 3;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+ }
+
+ #pragma acc loop worker
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ int x = i | j * 5;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-2.c
new file mode 100644
index 0000000..925f9a0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-2.c
@@ -0,0 +1,51 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared in a local scope, broadcasting
+ to vector-partitioned mode. Successive vector loops. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ int x = i ^ j * 3;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+
+ x = i | j * 5;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-3.c
new file mode 100644
index 0000000..6129523
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-3.c
@@ -0,0 +1,57 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared in a local scope, broadcasting
+ to vector-partitioned mode. Aggregate worker variable. */
+
+typedef struct
+{
+ int x, y;
+} vec2;
+
+int
+main (int argc, char* argv[])
+{
+ int i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ vec2 pt;
+
+ pt.x = i ^ j * 3;
+ pt.y = i | j * 5;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += pt.x * k;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += pt.y * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-4.c
new file mode 100644
index 0000000..4cec00e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-4.c
@@ -0,0 +1,60 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared in a local scope, broadcasting
+ to vector-partitioned mode. Addressable worker variable. */
+
+typedef struct
+{
+ int x, y;
+} vec2;
+
+int
+main (int argc, char* argv[])
+{
+ int i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ vec2 pt, *ptp;
+
+ ptp = &pt;
+
+ pt.x = i ^ j * 3;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += ptp->x * k;
+
+ ptp->y = i | j * 5;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += pt.y * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-5.c
new file mode 100644
index 0000000..efc2206
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-5.c
@@ -0,0 +1,53 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared in a local scope, broadcasting
+ to vector-partitioned mode. Array worker variable. */
+
+int
+main (int argc, char* argv[])
+{
+ int i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ int pt[2];
+
+ pt[0] = i ^ j * 3;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += pt[0] * k;
+
+ pt[1] = i | j * 5;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += pt[1] * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-1.c
new file mode 100644
index 0000000..504d11e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-1.c
@@ -0,0 +1,27 @@
+#include <assert.h>
+
+/* Test of gang-private variables declared on loop directive. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32];
+
+ for (i = 0; i < 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ {
+ #pragma acc loop gang private(x)
+ for (i = 0; i < 32; i++)
+ {
+ x = i * 2;
+ arr[i] += x;
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ assert (arr[i] == i * 3);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-2.c
new file mode 100644
index 0000000..9debf83
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-2.c
@@ -0,0 +1,33 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of gang-private variables declared on loop directive, with broadcasting
+ to partitioned workers. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32 * 32];
+
+ for (i = 0; i < 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ {
+ #pragma acc loop gang private(x)
+ for (i = 0; i < 32; i++)
+ {
+ x = i * 2;
+
+ #pragma acc loop worker
+ for (int j = 0; j < 32; j++)
+ arr[i * 32 + j] += x;
+ }
+ }
+
+ for (i = 0; i < 32 * 32; i++)
+ assert (arr[i] == i + (i / 32) * 2);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-3.c
new file mode 100644
index 0000000..0a77869
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-3.c
@@ -0,0 +1,31 @@
+#include <assert.h>
+
+/* Test of gang-private variables declared on loop directive, with broadcasting
+ to partitioned vectors. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32 * 32];
+
+ for (i = 0; i < 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ {
+ #pragma acc loop gang private(x)
+ for (i = 0; i < 32; i++)
+ {
+ x = i * 2;
+
+ #pragma acc loop vector
+ for (int j = 0; j < 32; j++)
+ arr[i * 32 + j] += x;
+ }
+ }
+
+ for (i = 0; i < 32 * 32; i++)
+ assert (arr[i] == i + (i / 32) * 2);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-4.c
new file mode 100644
index 0000000..24b0059
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-4.c
@@ -0,0 +1,37 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of gang-private addressable variable declared on loop directive, with
+ broadcasting to partitioned workers. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32 * 32];
+
+ for (i = 0; i < 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ {
+ #pragma acc loop gang private(x)
+ for (i = 0; i < 32; i++)
+ {
+ int *p = &x;
+
+ x = i * 2;
+
+ #pragma acc loop worker
+ for (int j = 0; j < 32; j++)
+ arr[i * 32 + j] += x;
+
+ (*p)--;
+ }
+ }
+
+ for (i = 0; i < 32 * 32; i++)
+ assert (arr[i] == i + (i / 32) * 2);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-5.c
new file mode 100644
index 0000000..b955303
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-5.c
@@ -0,0 +1,34 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of gang-private array variable declared on loop directive, with
+ broadcasting to partitioned workers. */
+
+int
+main (int argc, char* argv[])
+{
+ int x[8], i, arr[32 * 32];
+
+ for (i = 0; i < 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ {
+ #pragma acc loop gang private(x)
+ for (i = 0; i < 32; i++)
+ {
+ for (int j = 0; j < 8; j++)
+ x[j] = j * 2;
+
+ #pragma acc loop worker
+ for (int j = 0; j < 32; j++)
+ arr[i * 32 + j] += x[j % 8];
+ }
+ }
+
+ for (i = 0; i < 32 * 32; i++)
+ assert (arr[i] == i + (i % 8) * 2);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-6.c
new file mode 100644
index 0000000..0c17eaa
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-6.c
@@ -0,0 +1,42 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of gang-private aggregate variable declared on loop directive, with
+ broadcasting to partitioned workers. */
+
+typedef struct {
+ int x, y, z;
+ int attr[13];
+} vec3;
+
+int
+main (int argc, char* argv[])
+{
+ int i, arr[32 * 32];
+ vec3 pt;
+
+ for (i = 0; i < 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ {
+ #pragma acc loop gang private(pt)
+ for (i = 0; i < 32; i++)
+ {
+ pt.x = i;
+ pt.y = i * 2;
+ pt.z = i * 4;
+ pt.attr[5] = i * 6;
+
+ #pragma acc loop worker
+ for (int j = 0; j < 32; j++)
+ arr[i * 32 + j] += pt.x + pt.y + pt.z + pt.attr[5];
+ }
+ }
+
+ for (i = 0; i < 32 * 32; i++)
+ assert (arr[i] == i + (i / 32) * 13);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-vector-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-vector-1.c
new file mode 100644
index 0000000..b3c6ad3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-vector-1.c
@@ -0,0 +1,51 @@
+#include <assert.h>
+
+/* Test of vector-private variables declared on loop directive. */
+
+int
+main (int argc, char* argv[])
+{
+ int x, i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+
+ #pragma acc loop vector private(x)
+ for (k = 0; k < 32; k++)
+ {
+ x = i ^ j * 3;
+ arr[i * 1024 + j * 32 + k] += x * k;
+ }
+
+ #pragma acc loop vector private(x)
+ for (k = 0; k < 32; k++)
+ {
+ x = i | j * 5;
+ arr[i * 1024 + j * 32 + k] += x * k;
+ }
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-vector-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-vector-2.c
new file mode 100644
index 0000000..d4609e9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-vector-2.c
@@ -0,0 +1,46 @@
+#include <assert.h>
+
+/* Test of vector-private variables declared on loop directive. Array type. */
+
+int
+main (int argc, char* argv[])
+{
+ int pt[2], i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+
+ #pragma acc loop vector private(pt)
+ for (k = 0; k < 32; k++)
+ {
+ pt[0] = i ^ j * 3;
+ pt[1] = i | j * 5;
+ arr[i * 1024 + j * 32 + k] += pt[0] * k;
+ arr[i * 1024 + j * 32 + k] += pt[1] * k;
+ }
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-1.c
new file mode 100644
index 0000000..6377f89
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-1.c
@@ -0,0 +1,36 @@
+#include <assert.h>
+
+/* Test of worker-private variables declared on a loop directive. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32 * 32];
+
+ for (i = 0; i < 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(8) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker private(x)
+ for (j = 0; j < 32; j++)
+ {
+ x = i ^ j * 3;
+ /* Try to ensure 'x' accesses doesn't get optimized into a
+ temporary. */
+ __asm__ __volatile__ ("");
+ arr[i * 32 + j] += x;
+ }
+ }
+ }
+
+ for (i = 0; i < 32 * 32; i++)
+ assert (arr[i] == i + ((i / 32) ^ (i % 32) * 3));
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-2.c
new file mode 100644
index 0000000..834b0f1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-2.c
@@ -0,0 +1,45 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared on a loop directive, broadcasting
+ to vector-partitioned mode. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker private(x)
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ x = i ^ j * 3;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-3.c
new file mode 100644
index 0000000..18d6229
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-3.c
@@ -0,0 +1,56 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared on a loop directive, broadcasting
+ to vector-partitioned mode. Back-to-back worker loops. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker private(x)
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ x = i ^ j * 3;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+ }
+
+ #pragma acc loop worker private(x)
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ x = i | j * 5;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-4.c
new file mode 100644
index 0000000..881bc7d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-4.c
@@ -0,0 +1,51 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared on a loop directive, broadcasting
+ to vector-partitioned mode. Successive vector loops. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker private(x)
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ x = i ^ j * 3;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+
+ x = i | j * 5;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-5.c
new file mode 100644
index 0000000..fc1535b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-5.c
@@ -0,0 +1,53 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared on a loop directive, broadcasting
+ to vector-partitioned mode. Addressable worker variable. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32 * 32 * 32];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker private(x)
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+ int *p = &x;
+
+ x = i ^ j * 3;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+
+ *p = i | j * 5;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += x * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-6.c
new file mode 100644
index 0000000..feba09e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-6.c
@@ -0,0 +1,57 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared on a loop directive, broadcasting
+ to vector-partitioned mode. Aggregate worker variable. */
+
+typedef struct
+{
+ int x, y;
+} vec2;
+
+int
+main (int argc, char* argv[])
+{
+ int i, arr[32 * 32 * 32];
+ vec2 pt;
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ #pragma acc loop worker private(pt)
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+
+ pt.x = i ^ j * 3;
+ pt.y = i | j * 5;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += pt.x * k;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += pt.y * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-7.c
new file mode 100644
index 0000000..5469c5d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-7.c
@@ -0,0 +1,56 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of worker-private variables declared on loop directive, broadcasting
+ to vector-partitioned mode. Array worker variable. */
+
+int
+main (int argc, char* argv[])
+{
+ int i, arr[32 * 32 * 32];
+ int pt[2];
+
+ for (i = 0; i < 32 * 32 * 32; i++)
+ arr[i] = i;
+
+ /* "pt" is treated as "present_or_copy" on the parallel directive because it
+ is an array variable. */
+ #pragma acc parallel copy(arr) num_gangs(32) num_workers(32) vector_length(32)
+ {
+ int j;
+
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ /* But here, it is made private per-worker. */
+ #pragma acc loop worker private(pt)
+ for (j = 0; j < 32; j++)
+ {
+ int k;
+
+ pt[0] = i ^ j * 3;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += pt[0] * k;
+
+ pt[1] = i | j * 5;
+
+ #pragma acc loop vector
+ for (k = 0; k < 32; k++)
+ arr[i * 1024 + j * 32 + k] += pt[1] * k;
+ }
+ }
+ }
+
+ for (i = 0; i < 32; i++)
+ for (int j = 0; j < 32; j++)
+ for (int k = 0; k < 32; k++)
+ {
+ int idx = i * 1024 + j * 32 + k;
+ assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-1.c
new file mode 100644
index 0000000..831ba1e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-1.c
@@ -0,0 +1,25 @@
+#include <assert.h>
+
+/* Basic test of firstprivate variable. */
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[32];
+
+ for (i = 0; i < 32; i++)
+ arr[i] = 3;
+
+ #pragma acc parallel firstprivate(x) copy(arr) num_gangs(32) num_workers(8) \
+ vector_length(32)
+ {
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ arr[i] += x;
+ }
+
+ for (i = 0; i < 32; i++)
+ assert (arr[i] == 8);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-2.c
new file mode 100644
index 0000000..8199186
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-2.c
@@ -0,0 +1,35 @@
+#include <assert.h>
+
+/* Test of gang-private variables declared on the parallel directive. */
+
+#if defined(ACC_DEVICE_TYPE_host) || defined(ACC_DEVICE_TYPE_host_nonshm)
+#define ACTUAL_GANGS 1
+#else
+#define ACTUAL_GANGS 32
+#endif
+
+int
+main (int argc, char* argv[])
+{
+ int x = 5, i, arr[ACTUAL_GANGS];
+
+ for (i = 0; i < ACTUAL_GANGS; i++)
+ arr[i] = 3;
+
+ #pragma acc parallel private(x) copy(arr) num_gangs(ACTUAL_GANGS) \
+ num_workers(8) vector_length(32)
+ {
+ #pragma acc loop gang(static:1)
+ for (i = 0; i < ACTUAL_GANGS; i++)
+ x = i * 2;
+
+ #pragma acc loop gang(static:1)
+ for (i = 0; i < ACTUAL_GANGS; i++)
+ arr[i] += x;
+ }
+
+ for (i = 0; i < ACTUAL_GANGS; i++)
+ assert (arr[i] == 3 + i * 2);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-3.c
new file mode 100644
index 0000000..e5387c6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-3.c
@@ -0,0 +1,35 @@
+/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+
+#include <assert.h>
+
+/* Test of gang-private array variable declared on the parallel directive. */
+
+int
+main (int argc, char* argv[])
+{
+ int x[32], i, arr[32 * 32];
+
+ for (i = 0; i < 32 * 32; i++)
+ arr[i] = i;
+
+ #pragma acc parallel private(x) copy(arr) num_gangs(32) num_workers(2) \
+ vector_length(32)
+ {
+ #pragma acc loop gang
+ for (i = 0; i < 32; i++)
+ {
+ int j;
+ for (j = 0; j < 32; j++)
+ x[j] = j * 2;
+
+ #pragma acc loop worker
+ for (j = 0; j < 32; j++)
+ arr[i * 32 + j] += x[31 - j];
+ }
+ }
+
+ for (i = 0; i < 32 * 32; i++)
+ assert (arr[i] == i + (31 - (i % 32)) * 2);
+
+ return 0;
+}
^ permalink raw reply [flat|nested] 2+ messages in thread
* Re: [gomp4] Tests for private variables/state propagation
2015-06-17 14:23 [gomp4] Tests for private variables/state propagation Julian Brown
@ 2015-06-18 14:33 ` Thomas Schwinge
0 siblings, 0 replies; 2+ messages in thread
From: Thomas Schwinge @ 2015-06-18 14:33 UTC (permalink / raw)
To: Julian Brown, Cesar Philippidis, Bernd Schmidt; +Cc: gcc-patches, nathan
[-- Attachment #1: Type: text/plain, Size: 12074 bytes --]
Hi!
On Wed, 17 Jun 2015 15:15:15 +0100, Julian Brown <julian@codesourcery.com> wrote:
> This is a set of tests for OpenACC private variable/state propagation
> support in GCC. The associated functionality is a work-in-progress: as
> such, many of these tests do not pass yet (causing incorrect results,
> ICEs or even bogus assembly output). I believe the tests to be valid
> OpenACC, though it's possible I misinterpreted the spec at some points!
>
> I will apply to the gomp4 branch shortly. (We will of course be working
> on addressing the failures.)
Thanks for these! To avoid FAILs in the test logs, based on results of
running the testsuite on my laptop as well as a system with Nvidia K20
hardware, in r224614, I tuned the XFAILs as follows. Please shout if
this doesn't match your results.
I noticed one thing, and unless you (Julian, Bernd) are aware of that
already, would you (Cesar) please look into figuringing out why there are
several execution tests that FAIL for C but PASS for C++? (Maybe the
respective C and C++ tree dumps are similar enough to make a comparison
feasible?) If the FAILs can be attributed to some missing implementation
in OMP lowering/expansion, then that's fine (that is, currently expected)
-- but why then do the tests PASS for C++?
Curiously, there is one test case,
libgomp.oacc-c-c++-common/private-vars-loop-gang-3.c, that started
FAILing (again for C only; »private-vars-loop-gang-3.exe:
[...]/private-vars-loop-gang-3.c:28: main: Assertion `arr[i] == i + (i /
32) * 2' failed.«) only after your (Cesar's) recent gomp-4_0-branch,
r224586 »Only use static_chunk scheduling for acc gang loops when a the
gang clause contains a non-'*' static argument«. Maybe that can give a
clue, maybe not.
commit d29c4d788a5857884584f048c45f6959643b98eb
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Jun 18 14:14:40 2015 +0000
Tune XFAILs
... for tests added in gomp-4_0-branch, r224561.
libgomp/
* testsuite/libgomp.oacc-c++/c++.exp (check_effective_target_c):
New procedure.
* testsuite/libgomp.oacc-c/c.exp (check_effective_target_c):
Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-1.c:
Tune XFAILs.
* testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-5.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-3.c:
Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@224614 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog.gomp | 27 ++++++++++++++++++++
libgomp/testsuite/libgomp.oacc-c++/c++.exp | 4 +++
.../private-vars-local-worker-1.c | 2 +-
.../private-vars-local-worker-2.c | 2 +-
.../private-vars-local-worker-4.c | 1 +
.../private-vars-loop-gang-3.c | 6 +++++
.../private-vars-loop-gang-4.c | 3 ++-
.../private-vars-loop-worker-2.c | 2 +-
.../private-vars-loop-worker-3.c | 2 +-
.../private-vars-loop-worker-4.c | 2 +-
.../private-vars-loop-worker-5.c | 3 ++-
.../private-vars-par-gang-3.c | 2 +-
libgomp/testsuite/libgomp.oacc-c/c.exp | 4 +++
13 files changed, 52 insertions(+), 8 deletions(-)
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 8b8243a..5cd5d06 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,30 @@
+2015-06-18 Thomas Schwinge <thomas@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c++/c++.exp (check_effective_target_c):
+ New procedure.
+ * testsuite/libgomp.oacc-c/c.exp (check_effective_target_c):
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-1.c:
+ Tune XFAILs.
+ * testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-4.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-3.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-4.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-2.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-3.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-4.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-5.c:
+ Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-3.c:
+ Likewise.
+
2015-06-18 Tom de Vries <tom@codesourcery.com>
* testsuite/libgomp.fortran/parloops-exit-first-loop-alt-2.f95: New test.
diff --git libgomp/testsuite/libgomp.oacc-c++/c++.exp libgomp/testsuite/libgomp.oacc-c++/c++.exp
index 8d8f4ff..a3d0eda 100644
--- libgomp/testsuite/libgomp.oacc-c++/c++.exp
+++ libgomp/testsuite/libgomp.oacc-c++/c++.exp
@@ -14,6 +14,10 @@ if [info exists lang_include_flags] then {
unset lang_include_flags
}
+proc check_effective_target_c { } {
+ return 0
+}
+
# Initialize dg.
dg-init
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-1.c
index f8658e5..957f827 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-1.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-1.c
@@ -1,4 +1,4 @@
-/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+/* { dg-xfail-run-if "TODO" { c && openacc_nvidia_accel_selected } { "*" } { "" } } */
#include <assert.h>
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-2.c
index 925f9a0..5bf3e16 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-2.c
@@ -1,4 +1,4 @@
-/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+/* { dg-xfail-run-if "TODO" { c && openacc_nvidia_accel_selected } { "*" } { "" } } */
#include <assert.h>
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-4.c
index 4cec00e..118dae8 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-4.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-local-worker-4.c
@@ -1,3 +1,4 @@
+/* { dg-xfail-if "TODO" { c } } */
/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
#include <assert.h>
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-3.c
index 0a77869..bc96f0d 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-3.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-3.c
@@ -1,3 +1,9 @@
+/* { dg-xfail-run-if "TODO" { c && openacc_nvidia_accel_selected } { "*" } { "" } } */
+/* Fails as of gomp-4_0-branch r224586, but for C only:
+
+ private-vars-loop-gang-3.exe: [...]/private-vars-loop-gang-3.c:28: main: Assertion `arr[i] == i + (i / 32) * 2' failed.
+*/
+
#include <assert.h>
/* Test of gang-private variables declared on loop directive, with broadcasting
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-4.c
index 24b0059..f0f0477 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-4.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-gang-4.c
@@ -1,4 +1,5 @@
-/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+/* { dg-xfail-if "TODO: ICE" { *-*-* } } */
+/* { dg-excess-errors "TODO" } */
#include <assert.h>
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-2.c
index 834b0f1..a5f16d8 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-2.c
@@ -1,4 +1,4 @@
-/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+/* { dg-xfail-run-if "TODO" { c && openacc_nvidia_accel_selected } { "*" } { "" } } */
#include <assert.h>
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-3.c
index 18d6229..5c75f4d 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-3.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-3.c
@@ -1,4 +1,4 @@
-/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+/* { dg-xfail-run-if "TODO" { c && openacc_nvidia_accel_selected } { "*" } { "" } } */
#include <assert.h>
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-4.c
index 881bc7d..f0d1309 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-4.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-4.c
@@ -1,4 +1,4 @@
-/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+/* { dg-xfail-run-if "TODO" { c && openacc_nvidia_accel_selected } { "*" } { "" } } */
#include <assert.h>
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-5.c libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-5.c
index fc1535b..741795d 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-5.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-loop-worker-5.c
@@ -1,4 +1,5 @@
-/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+/* { dg-xfail-if "TODO: ICE" { *-*-* } } */
+/* { dg-excess-errors "TODO" } */
#include <assert.h>
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-3.c
index e5387c6..c580d68 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-3.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/private-vars-par-gang-3.c
@@ -1,4 +1,4 @@
-/* { dg-xfail-run-if "TODO" { openacc_nvidia_accel_selected } { "*" } { "" } } */
+/* { dg-xfail-run-if "TODO" { c && openacc_nvidia_accel_selected } { "*" } { "" } } */
#include <assert.h>
diff --git libgomp/testsuite/libgomp.oacc-c/c.exp libgomp/testsuite/libgomp.oacc-c/c.exp
index 8d54815..15e8f81 100644
--- libgomp/testsuite/libgomp.oacc-c/c.exp
+++ libgomp/testsuite/libgomp.oacc-c/c.exp
@@ -19,6 +19,10 @@ if ![info exists DEFAULT_CFLAGS] then {
set DEFAULT_CFLAGS "-O2"
}
+proc check_effective_target_c { } {
+ return 1
+}
+
# Initialize dg.
dg-init
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 2+ messages in thread
end of thread, other threads:[~2015-06-18 14:28 UTC | newest]
Thread overview: 2+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2015-06-17 14:23 [gomp4] Tests for private variables/state propagation Julian Brown
2015-06-18 14:33 ` Thomas Schwinge
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).