Tom de Vries 2b9d9e3937 [nvptx] Enable large vectors
Allow vector_length clauses to accept values larger than warp size.  Note that
this does not enable setting vector_length to values larger than warp size using
-fopenacc-dim.

2019-01-12  Tom de Vries  <tdevries@suse.de>

	* config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Take larger vector
	lengths into account.

	* testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Expect
	vector length to be 128.
	* testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Expect vector
	length 2097152 to be reduced to 1024 instead of 32.

From-SVN: r267889
2019-01-12 22:17:42 +00:00

39 lines
887 B
C

/* { 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-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow" } } */
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=128" } */