This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[og7, openacc, committed] Add vector-length-128-{1,2,3}.c test-cases
- From: Tom de Vries <Tom_deVries at mentor dot com>
- To: GCC Patches <gcc-patches at gcc dot gnu dot org>, Thomas Schwinge <thomas at codesourcery dot com>
- Cc: Cesar Philippidis <cesar_philippidis at mentor dot com>
- Date: Fri, 30 Mar 2018 15:31:40 +0200
- Subject: [og7, openacc, committed] Add vector-length-128-{1,2,3}.c test-cases
Hi,
this patch adds three testcases, setting vector length to 128 in three
different ways:
1. the vector_length clause
2. the -fopenacc-dim option
3. the GOMP_OPENACC_DIM variable
The tests contains:
- a check of the dimensions that the compiler decides upon
- a check of the dimensions used at runtime by libgomp
[ The first check is made possible by the "Add scan-offload-tree-dump"
patch I've just committed.
The second used setting environment variable GOMP_DEBUG to 1. We cannot
do this in main using setenv, as we can for GOMP_OPENACC_DIM, so we have
to use dg-set-target-env-var. This means the dg-output scans will fail
in a remote test setup where dg-set-target-env-var is broken. This is
annoying, but not as annoying as not having tests that DTRT in a local
test setup.
Also annoying is the fact that GOMP_DEBUG=1 prints a lot of lines into
libgomp.log. We could improve on that by not emitting the PTX code for
GOMP_DEBUG=1, and maybe emit that depending on the value of another
variable, say GOMP_OPENACC_DEBUG=n.
]
Currently, in all three test cases we don't use vector length 128, but
fall back to warp_size, so checks test for 32, not 128 vector_length.
Tested libgomp on x86_64 build with nvptx accelerator.
Committed.
Thanks,
- Tom
[openacc] Add vector-length-128-{1,2,3}.c test-cases
2018-03-30 Tom de Vries <tom@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: New test.
---
.../vector-length-128-1.c | 39 ++++++++++++++++++++
.../vector-length-128-2.c | 40 +++++++++++++++++++++
.../vector-length-128-3.c | 42 ++++++++++++++++++++++
3 files changed, 121 insertions(+)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c
new file mode 100644
index 0000000..fab5b0d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c
@@ -0,0 +1,39 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
+
+#include <stdlib.h>
+
+#define N 1024
+
+unsigned int a[N];
+unsigned int b[N];
+unsigned int c[N];
+unsigned int n = N;
+
+int
+main (void)
+{
+ for (unsigned int i = 0; i < n; ++i)
+ {
+ a[i] = i % 3;
+ b[i] = i % 5;
+ }
+
+#pragma acc parallel vector_length (128) copyin (a,b) copyout (c)
+ {
+#pragma acc loop vector
+ for (unsigned int i = 0; i < n; i++)
+ c[i] = a[i] + b[i];
+ }
+
+ for (unsigned int i = 0; i < n; ++i)
+ if (c[i] != (i % 3) + (i % 5))
+ abort ();
+
+ return 0;
+}
+/* { dg-prune-output "using vector_length \\(32\\), ignoring 128" } */
+
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c
new file mode 100644
index 0000000..cc6fd55
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c
@@ -0,0 +1,40 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-fopenacc-dim=-:-:128" } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
+
+#include <stdlib.h>
+
+#define N 1024
+
+unsigned int a[N];
+unsigned int b[N];
+unsigned int c[N];
+unsigned int n = N;
+
+int
+main (void)
+{
+ for (unsigned int i = 0; i < n; ++i)
+ {
+ a[i] = i % 3;
+ b[i] = i % 5;
+ }
+
+#pragma acc parallel copyin (a,b) copyout (c)
+ {
+#pragma acc loop vector
+ for (unsigned int i = 0; i < n; i++)
+ c[i] = a[i] + b[i];
+ }
+
+ for (unsigned int i = 0; i < n; ++i)
+ if (c[i] != (i % 3) + (i % 5))
+ abort ();
+
+ return 0;
+}
+/* { dg-prune-output "using vector_length \\(32\\), ignoring 128" } */
+
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c
new file mode 100644
index 0000000..c403e74
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c
@@ -0,0 +1,42 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* We default to warp size 32 for the vector length, so the GOMP_OPENACC_DIM has
+ no effect. */
+/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "-:-:128" } */
+/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
+
+
+#include <stdlib.h>
+
+#define N 1024
+
+unsigned int a[N];
+unsigned int b[N];
+unsigned int c[N];
+unsigned int n = N;
+
+int
+main (void)
+{
+ for (unsigned int i = 0; i < n; ++i)
+ {
+ a[i] = i % 3;
+ b[i] = i % 5;
+ }
+
+#pragma acc parallel copyin (a,b) copyout (c)
+ {
+#pragma acc loop vector
+ for (unsigned int i = 0; i < n; i++)
+ c[i] = a[i] + b[i];
+ }
+
+ for (unsigned int i = 0; i < n; ++i)
+ if (c[i] != (i % 3) + (i % 5))
+ abort ();
+
+ return 0;
+}
+
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */