On 14-12-18 20:58, Tom de Vries wrote:
> 0016-nvptx-Add-vector_length-128-testcases.patch
> * testsuite/libgomp.oacc-c-c++-common/vred2d-128.c: New test.
>
> * testsuite/libgomp.oacc-fortran/gemm.f90: New test.
>
> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c: New
> test.
These test-cases fail at the point that they're introduced in the patch
series, which is not a good idea.
> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: New
> test.
> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: New
> test.
These test-cases could have been committed as obvious.
Committed last two as attached.
Thanks,
- Tom
[nvptx] Add vector_length 128 testcases
Add a couple of test-cases using vector length 128, while checking that we
override to vector length 32.
2018-12-17 Tom de Vries <tdevr...@suse.de>
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.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-3.c | 42 ++++++++++++++++++++++
2 files changed, 81 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 00000000000..fab5b0d25d1
--- /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-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c
new file mode 100644
index 00000000000..c403e74658b
--- /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" } */
