7ec16b79f0
2017-12-27 Tom de Vries <tom@codesourcery.com> PR c++/83046 * testsuite/libgomp.oacc-c-c++-common/gang-static-2.c (test_static) (test_nonstatic): Fix return type to workaround PR83046. From-SVN: r256008
101 lines
2.1 KiB
C
101 lines
2.1 KiB
C
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
|
/* This code uses nvptx inline assembly guarded with acc_on_device, which is
|
|
not optimized away at -O0, and then confuses the target assembler.
|
|
{ dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
|
|
|
|
#include <assert.h>
|
|
#include <openacc.h>
|
|
|
|
#define N 100
|
|
|
|
#define GANG_ID(I) \
|
|
(acc_on_device (acc_device_nvidia) \
|
|
? ({unsigned __r; \
|
|
__asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (__r)); \
|
|
__r; }) : (I))
|
|
|
|
void
|
|
test_static(int *a, int num_gangs, int sarg)
|
|
{
|
|
int i, j;
|
|
|
|
if (sarg == 0)
|
|
sarg = 1;
|
|
|
|
for (i = 0; i < N / sarg; i++)
|
|
for (j = 0; j < sarg; j++)
|
|
assert (a[i*sarg+j] == i % num_gangs);
|
|
}
|
|
|
|
void
|
|
test_nonstatic(int *a, int gangs)
|
|
{
|
|
int i, j;
|
|
|
|
for (i = 0; i < N; i+=gangs)
|
|
for (j = 0; j < gangs; j++)
|
|
assert (a[i+j] == i/gangs);
|
|
}
|
|
|
|
int
|
|
main ()
|
|
{
|
|
int a[N];
|
|
int i, x;
|
|
|
|
#pragma acc parallel loop gang (static:*) num_gangs (10)
|
|
for (i = 0; i < 100; i++)
|
|
a[i] = GANG_ID (i);
|
|
|
|
test_nonstatic (a, 10);
|
|
|
|
#pragma acc parallel loop gang (static:1) num_gangs (10)
|
|
for (i = 0; i < 100; i++)
|
|
a[i] = GANG_ID (i);
|
|
|
|
test_static (a, 10, 1);
|
|
|
|
#pragma acc parallel loop gang (static:2) num_gangs (10)
|
|
for (i = 0; i < 100; i++)
|
|
a[i] = GANG_ID (i);
|
|
|
|
test_static (a, 10, 2);
|
|
|
|
#pragma acc parallel loop gang (static:5) num_gangs (10)
|
|
for (i = 0; i < 100; i++)
|
|
a[i] = GANG_ID (i);
|
|
|
|
test_static (a, 10, 5);
|
|
|
|
#pragma acc parallel loop gang (static:20) num_gangs (10)
|
|
for (i = 0; i < 100; i++)
|
|
a[i] = GANG_ID (i);
|
|
|
|
test_static (a, 10, 20);
|
|
|
|
/* Non-static gang. */
|
|
#pragma acc parallel loop gang num_gangs (10)
|
|
for (i = 0; i < 100; i++)
|
|
a[i] = GANG_ID (i);
|
|
|
|
test_nonstatic (a, 10);
|
|
|
|
/* Static arguments with a variable expression. */
|
|
|
|
x = 20;
|
|
#pragma acc parallel loop gang (static:0+x) num_gangs (10)
|
|
for (i = 0; i < 100; i++)
|
|
a[i] = GANG_ID (i);
|
|
|
|
test_static (a, 10, 20);
|
|
|
|
x = 20;
|
|
#pragma acc parallel loop gang (static:x) num_gangs (10)
|
|
for (i = 0; i < 100; i++)
|
|
a[i] = GANG_ID (i);
|
|
|
|
test_static (a, 10, 20);
|
|
|
|
return 0;
|
|
}
|