mirrors/gcc
0
0
Fork 0
mirror of https://gcc.gnu.org/git/gcc.git synced 2024-12-01 07:44:05 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

[nvptx] Add vector_length 128 testcases

Browse Source 
Add a couple of test-cases using vector length 128, while checking that we
override to vector length 32.

2019-01-03  Tom de Vries  <tdevries@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.

From-SVN: r267559
This commit is contained in:
Tom de Vries 2019-01-03 15:08:46 +00:00 committed by Tom de Vries
parent 0d06e8c794
commit 5c571497e1
3 changed files with 86 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
libgomp/ChangeLog
View File

@ -1,3 +1,8 @@
2019-01-03 Tom de Vries <tdevries@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.
2019-01-01 Jakub Jelinek <jakub@redhat.com>
Update copyright years.

39
libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c Normal file
View File

@ -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" } */

42
libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c Normal file
View File

@ -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" } */