diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 7d054c27a91..021d01b7ff0 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2019-01-03 Tom de Vries + + * 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. + 2019-01-01 Jakub Jelinek Update copyright years. 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 + +#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 + +#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" } */