b97e78b712
gcc/ * builtins.c (expand_builtin_acc_on_device) [ACCEL_COMPILER]: Emit open-coded sequence. * omp-low.c (oacc_process_reduction_data): Remove handline of GOMP_DEVICE_HOST_NONSHM. gcc/testsuite/ * c-c++-common/goacc/acc_on_device-2.c: Remove XFAIL for C. include/ * gomp-constants.c (GOMP_DEVICE_HOST_NONSHM): Remove. libgomp/ * libgomp-plugin.h (enum offload_target_type): Remove OFFLOAD_TARGET_TYPE_HOST_NONSHM. * openacc.f90 (openacc_kinds): Remove acc_device_host_nonshm. * openacc.h (enum acc_device_t): Likewise. * openacc_lib.h: Likewise. * oacc-init.c (name_of_acc_device_t): Don't handle it. (acc_on_device): Just use __builtin_acc_on_device. * testsuite/libgomp.oacc-c-c++-common/if-1.c: Don't forbid usage of acc_on_device builtin. * plugin/plugin-host.h: Remove file. * plugin/plugin-host.c: Likewise, but salvage some content into... * oacc-host.c: ... this file. * plugin/Makefrag.am: Don't build libgomp-plugin-host_nonshm.la. * plugin/configfrag.ac (offload_targets): Don't add host_nonshm. * Makefile.in: Regenerate. * configure: Likewise. * testsuite/lib/libgomp.exp (check_effective_target_openacc_host_nonshm_selected): Remove. * testsuite/libgomp.oacc-c++/c++.exp: Don't handle ACC_DEVICE_TYPE=host_nonshm. * testsuite/libgomp.oacc-c/c.exp: Likewise. * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_on_device-1.c: Likewise. * testsuite/libgomp.oacc-fortran/acc_on_device-1-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/acc_on_device-1-2.f: Likewise. * testsuite/libgomp.oacc-fortran/acc_on_device-1-3.f: Likewise. From-SVN: r226763
613 lines
10 KiB
C
613 lines
10 KiB
C
/* { dg-do run } */
|
|
|
|
#include <openacc.h>
|
|
#include <stdlib.h>
|
|
#include <stdbool.h>
|
|
|
|
#define N 32
|
|
|
|
int
|
|
main(int argc, char **argv)
|
|
{
|
|
float *a, *b, *d_a, *d_b, exp, exp2;
|
|
int i;
|
|
const int one = 1;
|
|
const int zero = 0;
|
|
int n;
|
|
|
|
a = (float *) malloc (N * sizeof (float));
|
|
b = (float *) malloc (N * sizeof (float));
|
|
d_a = (float *) acc_malloc (N * sizeof (float));
|
|
d_b = (float *) acc_malloc (N * sizeof (float));
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 4.0;
|
|
|
|
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
#if ACC_MEM_SHARED
|
|
exp = 5.0;
|
|
#else
|
|
exp = 4.0;
|
|
#endif
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != exp)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 16.0;
|
|
|
|
#pragma acc parallel if(0)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != 17.0)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 8.0;
|
|
|
|
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(one)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
#if ACC_MEM_SHARED
|
|
exp = 9.0;
|
|
#else
|
|
exp = 8.0;
|
|
#endif
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != exp)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 22.0;
|
|
|
|
#pragma acc parallel if(zero)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != 23.0)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 16.0;
|
|
|
|
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(true)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
#if ACC_MEM_SHARED
|
|
exp = 17.0;
|
|
#else
|
|
exp = 16.0;
|
|
#endif
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != exp)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 76.0;
|
|
|
|
#pragma acc parallel if(false)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != 77.0)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 22.0;
|
|
|
|
n = 1;
|
|
|
|
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(n)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
#if ACC_MEM_SHARED
|
|
exp = 23.0;
|
|
#else
|
|
exp = 22.0;
|
|
#endif
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != exp)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 18.0;
|
|
|
|
n = 0;
|
|
|
|
#pragma acc parallel if(n)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != 19.0)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 49.0;
|
|
|
|
n = 1;
|
|
|
|
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(n + n)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
#if ACC_MEM_SHARED
|
|
exp = 50.0;
|
|
#else
|
|
exp = 49.0;
|
|
#endif
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != exp)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 38.0;
|
|
|
|
n = 0;
|
|
|
|
#pragma acc parallel if(n + n)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != 39.0)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 91.0;
|
|
|
|
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(-2)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
#if ACC_MEM_SHARED
|
|
exp = 92.0;
|
|
#else
|
|
exp = 91.0;
|
|
#endif
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != exp)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 43.0;
|
|
|
|
#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(one == 1)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
#if ACC_MEM_SHARED
|
|
exp = 44.0;
|
|
#else
|
|
exp = 43.0;
|
|
#endif
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != exp)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 87.0;
|
|
|
|
#pragma acc parallel if(one == 0)
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
if (acc_on_device (acc_device_host))
|
|
b[ii] = a[ii] + 1;
|
|
else
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != 88.0)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 3.0;
|
|
b[i] = 9.0;
|
|
}
|
|
|
|
#if ACC_MEM_SHARED
|
|
exp = 0.0;
|
|
exp2 = 0.0;
|
|
#else
|
|
acc_map_data (a, d_a, N * sizeof (float));
|
|
acc_map_data (b, d_b, N * sizeof (float));
|
|
exp = 3.0;
|
|
exp2 = 9.0;
|
|
#endif
|
|
|
|
#pragma acc update device(a[0:N], b[0:N]) if(1)
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 0.0;
|
|
b[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc update host(a[0:N], b[0:N]) if(1)
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != exp)
|
|
abort();
|
|
|
|
if (b[i] != exp2)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 6.0;
|
|
b[i] = 12.0;
|
|
}
|
|
|
|
#pragma acc update device(a[0:N], b[0:N]) if(0)
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 0.0;
|
|
b[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc update host(a[0:N], b[0:N]) if(1)
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != exp)
|
|
abort();
|
|
|
|
if (b[i] != exp2)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 26.0;
|
|
b[i] = 21.0;
|
|
}
|
|
|
|
#pragma acc update device(a[0:N], b[0:N]) if(1)
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 0.0;
|
|
b[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc update host(a[0:N], b[0:N]) if(0)
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (a[i] != 0.0)
|
|
abort();
|
|
|
|
if (b[i] != 0.0)
|
|
abort();
|
|
}
|
|
|
|
#if !ACC_MEM_SHARED
|
|
acc_unmap_data (a);
|
|
acc_unmap_data (b);
|
|
#endif
|
|
|
|
acc_free (d_a);
|
|
acc_free (d_b);
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 4.0;
|
|
b[i] = 0.0;
|
|
}
|
|
|
|
#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(1)
|
|
{
|
|
#pragma acc parallel present(a[0:N])
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != 4.0)
|
|
abort();
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 8.0;
|
|
b[i] = 1.0;
|
|
}
|
|
|
|
#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(0)
|
|
{
|
|
#if !ACC_MEM_SHARED
|
|
if (acc_is_present (a, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
a[i] = 18.0;
|
|
b[i] = 21.0;
|
|
}
|
|
|
|
#pragma acc data copyin(a[0:N]) if(1)
|
|
{
|
|
#if !ACC_MEM_SHARED
|
|
if (!acc_is_present (a, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#pragma acc data copyout(b[0:N]) if(0)
|
|
{
|
|
#if !ACC_MEM_SHARED
|
|
if (acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#pragma acc data copyout(b[0:N]) if(1)
|
|
{
|
|
#pragma acc parallel present(a[0:N]) present(b[0:N])
|
|
{
|
|
int ii;
|
|
|
|
for (ii = 0; ii < N; ii++)
|
|
{
|
|
b[ii] = a[ii];
|
|
}
|
|
}
|
|
}
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
if (b[i] != 18.0)
|
|
abort ();
|
|
}
|
|
|
|
#pragma acc enter data copyin (b[0:N]) if (0)
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#pragma acc exit data delete (b[0:N]) if (0)
|
|
|
|
#pragma acc enter data copyin (b[0:N]) if (1)
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (!acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#pragma acc exit data delete (b[0:N]) if (1)
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#pragma acc enter data copyin (b[0:N]) if (zero)
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#pragma acc exit data delete (b[0:N]) if (zero)
|
|
|
|
#pragma acc enter data copyin (b[0:N]) if (one)
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (!acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#pragma acc exit data delete (b[0:N]) if (one)
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#pragma acc enter data copyin (b[0:N]) if (one == 0)
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#pragma acc exit data delete (b[0:N]) if (one == 0)
|
|
|
|
#pragma acc enter data copyin (b[0:N]) if (one == 1)
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (!acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
#pragma acc exit data delete (b[0:N]) if (one == 1)
|
|
|
|
#if !ACC_MEM_SHARED
|
|
if (acc_is_present (b, N * sizeof (float)))
|
|
abort ();
|
|
#endif
|
|
|
|
return 0;
|
|
}
|