2620c80db0
gcc/testsuite/ * c-c++-common/goacc/combined-directives.c: Clean up dg-* directives. * c-c++-common/goacc/loop-clauses.c: Likewise. * g++.dg/goacc/template.C: Likewise. * gfortran.dg/goacc/combined-directives.f90: Likewise. * gfortran.dg/goacc/loop-1.f95: Likewise. * gfortran.dg/goacc/loop-5.f95: Likewise. * gfortran.dg/goacc/loop-6.f95: Likewise. * gfortran.dg/goacc/loop-tree-1.f90: Likewise. * c-c++-common/goacc-gomp/nesting-1.c: Update. * c-c++-common/goacc-gomp/nesting-fail-1.c: Likewise. * c-c++-common/goacc/clauses-fail.c: Likewise. * c-c++-common/goacc/parallel-1.c: Likewise. * c-c++-common/goacc/reduction-1.c: Likewise. * c-c++-common/goacc/reduction-2.c: Likewise. * c-c++-common/goacc/reduction-3.c: Likewise. * c-c++-common/goacc/reduction-4.c: Likewise. * c-c++-common/goacc/routine-3.c: Likewise. * c-c++-common/goacc/routine-4.c: Likewise. * c-c++-common/goacc/routine-5.c: Likewise. * c-c++-common/goacc/tile.c: Likewise. * g++.dg/goacc/template.C: Likewise. * gfortran.dg/goacc/combined-directives.f90: Likewise. * c-c++-common/goacc/nesting-1.c: Move dg-error test cases into... * c-c++-common/goacc/nesting-fail-1.c: ... this file. Update. * c-c++-common/goacc/kernels-1.c: Update. Incorporate... * c-c++-common/goacc/kernels-empty.c: ... this file, and... * c-c++-common/goacc/kernels-eternal.c: ... this file, and... * c-c++-common/goacc/kernels-noreturn.c: ... this file. * c-c++-common/goacc/host_data-1.c: New file. Incorporate... * c-c++-common/goacc/use_device-1.c: ... this file. * c-c++-common/goacc/host_data-2.c: New file. Incorporate... * c-c++-common/goacc/host_data-5.c: ... this file, and... * c-c++-common/goacc/host_data-6.c: ... this file. * c-c++-common/goacc/loop-2-kernels.c: New file. * c-c++-common/goacc/loop-2-parallel.c: Likewise. * c-c++-common/goacc/loop-3.c: Likewise. * g++.dg/goacc/reference.C: Likewise. * g++.dg/goacc/routine-1.C: Likewise. * g++.dg/goacc/routine-2.C: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/deviceptr-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/if-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-loop.c: Likewise. * testsuite/libgomp.oacc-fortran/asyncwait-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/asyncwait-2.f90: Likewise. * testsuite/libgomp.oacc-fortran/asyncwait-3.f90: Likewise. * testsuite/libgomp.oacc-fortran/declare-1.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Likewise. XFAIL. * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Update. Incorporate... * testsuite/libgomp.oacc-c-c++-common/firstprivate-2.c: ... this file. * testsuite/libgomp.oacc-c++/template-reduction.C: New file. * testsuite/libgomp.oacc-c-c++-common/gang-static-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/gang-static-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-clauses.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/private-variables.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c: Likewise. * testsuite/libgomp.oacc-fortran/clauses-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/default-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/firstprivate-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/gang-static-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/if-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/implicit-firstprivate-ref.f90: Likewise. * testsuite/libgomp.oacc-fortran/pr68813.f90: Likewise. * testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-1.c: Merge this file... * testsuite/libgomp.oacc-c-c++-common/parallel-1.c: ..., and this file into... * testsuite/libgomp.oacc-c-c++-common/data-clauses.h: ... this new file. Update. * testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels.c: New file. * testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-2.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/data-clauses-kernels-ipa-pta.c: ... this new file. Update. * testsuite/libgomp.oacc-c-c++-common/parallel-2.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/data-clauses-parallel-ipa-pta.c: ... this new file. Update. * testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: New file. Incorporate... * testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/worker-single-4.c: ... this file, and... * testsuite/libgomp.oacc-c-c++-common/worker-single-6.c: ... this file. * testsuite/libgomp.oacc-c-c++-common/update-1-2.c: Remove file. Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com> Co-Authored-By: Chung-Lin Tang <cltang@codesourcery.com> Co-Authored-By: James Norris <jnorris@codesourcery.com> Co-Authored-By: Julian Brown <julian@codesourcery.com> Co-Authored-By: Nathan Sidwell <nathan@codesourcery.com> Co-Authored-By: Tom de Vries <tom@codesourcery.com> From-SVN: r234575
963 lines
16 KiB
C
963 lines
16 KiB
C
#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
|
|
|
|
for (i = 0; i < N; i++)
|
|
a[i] = 4.0;
|
|
|
|
#pragma acc kernels 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 kernels 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 kernels 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 kernels 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 kernels 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 kernels 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 kernels 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 kernels 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 kernels 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 kernels 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 kernels 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 kernels 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 kernels 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
|
|
|
|
return 0;
|
|
}
|