[nvptx] Enable large vectors -- test-cases
Add various test-cases with vector length 128. 2019-01-12 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: New test. From-SVN: r267891
This commit is contained in:
parent
52d22ece49
commit
8e77f71eda
|
@ -1,3 +1,9 @@
|
|||
2019-01-12 Tom de Vries <tdevries@suse.de>
|
||||
|
||||
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: New test.
|
||||
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: New test.
|
||||
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: New test.
|
||||
|
||||
2019-01-12 Tom de Vries <tdevries@suse.de>
|
||||
|
||||
* plugin/plugin-nvptx.c (nvptx_exec): Update insufficient hardware
|
||||
|
|
|
@ -0,0 +1,40 @@
|
|||
/* { 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 num_workers (2) vector_length (128) copyin (a,b) copyout (c)
|
||||
{
|
||||
#pragma acc loop worker
|
||||
for (unsigned int i = 0; i < 4; i++)
|
||||
#pragma acc loop vector
|
||||
for (unsigned int j = 0; j < n / 4; j++)
|
||||
c[(i * N / 4) + j] = a[(i * N / 4) + j] + b[(i * N / 4) + j];
|
||||
}
|
||||
|
||||
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, 2, 128\\)" "oaccdevlow" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */
|
|
@ -0,0 +1,41 @@
|
|||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" ":2:" } */
|
||||
/* { 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 worker
|
||||
for (unsigned int i = 0; i < 4; i++)
|
||||
#pragma acc loop vector
|
||||
for (unsigned int j = 0; j < n / 4; j++)
|
||||
c[(i * N / 4) + j] = a[(i * N / 4) + j] + b[(i * N / 4) + j];
|
||||
}
|
||||
|
||||
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, 0, 128\\)" "oaccdevlow" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */
|
|
@ -0,0 +1,40 @@
|
|||
/* { 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 worker
|
||||
for (unsigned int i = 0; i < 4; i++)
|
||||
#pragma acc loop vector
|
||||
for (unsigned int j = 0; j < n / 4; j++)
|
||||
c[(i * N / 4) + j] = a[(i * N / 4) + j] + b[(i * N / 4) + j];
|
||||
}
|
||||
|
||||
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, 0, 128\\)" "oaccdevlow" } } */
|
||||
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=8, vectors=128" } */
|
Loading…
Reference in New Issue