procyberian/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

[nvptx] Force vl32 if calling vector-partitionable routines -- test-cases

Browse source 
Add test-cases for "[nvptx] Force vl32 if calling vector-partitionable
routines".

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

	PR target/85486
	* testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/pr85486.c: New test.

From-SVN: r267894
This commit is contained in:
Tom de Vries 2019-01-12 22:18:50 +00:00 • committed by Tom de Vries
parent b39e4366a2
commit 56314b772f
3 changed files with 111 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

6
libgomp/ChangeLog
View file

@ -1,3 +1,9 @@
2019-01-12 Tom de Vries <tdevries@suse.de>
PR target/85486
* testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/pr85486.c: New test.
2019-01-12 Tom de Vries <tdevries@suse.de>
PR target/85381

54
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-3.c Normal file
View file

@ -0,0 +1,54 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "::128" } */
/* Minimized from ref-1.C. */
#include <stdio.h>
#pragma acc routine vector
void __attribute__((noinline, noclone))
Vector (int *ptr, int n, const int inc)
{
#pragma acc loop vector
for (unsigned ix = 0; ix < n; ix++)
ptr[ix] += inc;
}
int
main (void)
{
const int n = 32, m=32;
int ary[m][n];
unsigned ix, iy;
for (ix = m; ix--;)
for (iy = n; iy--;)
ary[ix][iy] = (1 << 16) + (ix << 8) + iy;
int err = 0;
#pragma acc parallel copy (ary)
{
Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
}
for (ix = m; ix--;)
for (iy = n; iy--;)
if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy))
{
printf ("ary[%u][%u] = %x expected %x\n",
ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy));
err++;
}
if (err)
{
printf ("%d failed\n", err);
return 1;
}
return 0;
}
/* { dg-prune-output "using vector_length \\(32\\), ignoring runtime setting" } */

51
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486.c Normal file
View file

@ -0,0 +1,51 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* Minimized from ref-1.C. */
#include <stdio.h>
#pragma acc routine vector
void __attribute__((noinline, noclone))
Vector (int *ptr, int n, const int inc)
{
#pragma acc loop vector
for (unsigned ix = 0; ix < n; ix++)
ptr[ix] += inc;
}
int
main (void)
{
const int n = 32, m=32;
int ary[m][n];
unsigned ix, iy;
for (ix = m; ix--;)
for (iy = n; iy--;)
ary[ix][iy] = (1 << 16) + (ix << 8) + iy;
int err = 0;
#pragma acc parallel copy (ary) vector_length (128) /* { dg-warning "using vector_length \\(32\\) due to call to vector-partitionable routine, ignoring 128" } */
{
Vector (&ary[0][0], m * n, (1 << 24) - (1 << 16));
}
for (ix = m; ix--;)
for (iy = n; iy--;)
if (ary[ix][iy] != ((1 << 24) + (ix << 8) + iy))
{
printf ("ary[%u][%u] = %x expected %x\n",
ix, iy, ary[ix][iy], ((1 << 24) + (ix << 8) + iy));
err++;
}
if (err)
{
printf ("%d failed\n", err);
return 1;
}
return 0;
}