4bf9e5a8a2
gcc/c-family/ * c-pragma.c (oacc_pragmas): Add "atomic". * c-pragma.h (pragma_kind): Add PRAGMA_OACC_ATOMIC. gcc/c/ * c-parser.c (c_parser_omp_construct): Handle PRAGMA_OACC_ATOMIC. gcc/cp/ * parser.c (cp_parser_omp_construct, cp_parser_pragma): Handle PRAGMA_OACC_ATOMIC. gcc/fortran/ * gfortran.h (gfc_statement): Add ST_OACC_ATOMIC, ST_OACC_END_ATOMIC. (gfc_exec_op): Add EXEC_OACC_ATOMIC. * match.h (gfc_match_oacc_atomic): New prototype. * openmp.c (gfc_match_omp_atomic, gfc_match_oacc_atomic): New wrapper functions around... (gfc_match_omp_oacc_atomic): ... this new function. (oacc_code_to_statement, gfc_resolve_oacc_directive): Handle EXEC_OACC_ATOMIC. * parse.c (decode_oacc_directive): Handle "atomic", "end atomic". (case_exec_markers): Add ST_OACC_ATOMIC. (gfc_ascii_statement): Handle ST_OACC_ATOMIC, ST_OACC_END_ATOMIC. (parse_omp_atomic): Rename to... (parse_omp_oacc_atomic): ... this new function. Add omp_p formal parameter. Adjust all users. (parse_executable): Handle ST_OACC_ATOMIC. (is_oacc): Handle EXEC_OACC_ATOMIC. * resolve.c (gfc_resolve_blocks, gfc_resolve_code): Handle EXEC_OACC_ATOMIC. * st.c (gfc_free_statement): Handle EXEC_OACC_ATOMIC. * trans-openmp.c (gfc_trans_oacc_directive): Handle EXEC_OACC_ATOMIC. * trans.c (trans_code): Handle EXEC_OACC_ATOMIC. gcc/ * builtins.def (DEF_GOMP_BUILTIN): Enable for flag_openacc. * omp-low.c (check_omp_nesting_restrictions): Allow GIMPLE_OMP_ATOMIC_LOAD, GIMPLE_OMP_ATOMIC_STORE inside OpenACC contexts. gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-fail-1.c: Move "atomic" tests from here to... * c-c++-common/goacc-gomp/nesting-1.c: ... here, and expect them to succeed. libgomp/ * testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/atomic_rw-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c: Likewise. * testsuite/libgomp.oacc-fortran/atomic_capture-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/atomic_rw-1.f90: New file. * testsuite/libgomp.oacc-fortran/atomic_update-1.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/worker-single-1a.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/worker-single-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/worker-single-6.c: Likewise. From-SVN: r229703
1627 lines
26 KiB
C
1627 lines
26 KiB
C
/* { dg-do run } */
|
|
|
|
#include <stdlib.h>
|
|
|
|
int
|
|
main(int argc, char **argv)
|
|
{
|
|
int iexp, igot, imax, imin;
|
|
long long lexp, lgot;
|
|
int N = 32;
|
|
int i;
|
|
int idata[N];
|
|
long long ldata[N];
|
|
float fexp, fgot;
|
|
float fdata[N];
|
|
|
|
igot = 1234;
|
|
iexp = 31;
|
|
|
|
for (i = 0; i < N; i++)
|
|
idata[i] = i;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; igot = i; }
|
|
}
|
|
|
|
imax = 0;
|
|
imin = N;
|
|
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
imax = idata[i] > imax ? idata[i] : imax;
|
|
imin = idata[i] < imin ? idata[i] : imin;
|
|
}
|
|
|
|
if (imax != 1234 || imin != 0)
|
|
abort ();
|
|
|
|
return 0;
|
|
|
|
igot = 0;
|
|
iexp = 32;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; igot++; }
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 0;
|
|
iexp = 32;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; ++igot; }
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 0;
|
|
iexp = 32;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
#pragma acc atomic capture
|
|
{ ++igot; idata[i] = igot; }
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 0;
|
|
iexp = 32;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
#pragma acc atomic capture
|
|
{ igot++; idata[i] = igot; }
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 32;
|
|
iexp = 0;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; igot--; }
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 32;
|
|
iexp = 0;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; --igot; }
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 32;
|
|
iexp = 0;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
#pragma acc atomic capture
|
|
{ --igot; idata[i] = igot; }
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 32;
|
|
iexp = 0;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
#pragma acc atomic capture
|
|
{ igot--; idata[i] = igot; }
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
/* BINOP = + */
|
|
igot = 0;
|
|
iexp = 32;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; igot += expr; }
|
|
}
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 0;
|
|
iexp = 32;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ igot += expr; idata[i] = igot; }
|
|
}
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 0;
|
|
iexp = 32;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; igot = igot + expr; }
|
|
}
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 0;
|
|
iexp = 32;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; igot = expr + igot; }
|
|
}
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 0;
|
|
iexp = 32;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ igot = igot + expr; idata[i] = igot; }
|
|
}
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
|
|
igot = 0;
|
|
iexp = 32;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ igot = expr + igot; idata[i] = igot; }
|
|
}
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
/* BINOP = * */
|
|
lgot = 1LL;
|
|
lexp = 1LL << 32;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot *= expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 1LL << 32;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot *= expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 1LL << 32;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = lgot * expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 1LL << 32;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = expr * lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 1LL << 32;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = lgot * expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 1LL << 32;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = expr * lgot; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
/* BINOP = - */
|
|
igot = 32;
|
|
iexp = 0;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; igot -= expr; }
|
|
}
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 32;
|
|
iexp = 0;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ igot -= expr; idata[i] = igot; }
|
|
}
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 32;
|
|
iexp = 0;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; igot = igot - expr; }
|
|
}
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 1;
|
|
iexp = 1;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ idata[i] = igot; igot = expr - igot; }
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
if (i % 2 == 0)
|
|
{
|
|
if (idata[i] != 1)
|
|
abort ();
|
|
}
|
|
else
|
|
{
|
|
if (idata[i] != 0)
|
|
abort ();
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 1;
|
|
iexp = -31;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ igot = igot - expr; idata[i] = igot; }
|
|
}
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
igot = 1;
|
|
iexp = 1;
|
|
|
|
#pragma acc data copy (igot, idata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
int expr = 1;
|
|
|
|
#pragma acc atomic capture
|
|
{ igot = expr - igot; idata[i] = igot; }
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
if (i % 2 == 0)
|
|
{
|
|
if (idata[i] != 0)
|
|
abort ();
|
|
}
|
|
else
|
|
{
|
|
if (idata[i] != 1)
|
|
abort ();
|
|
}
|
|
|
|
if (iexp != igot)
|
|
abort ();
|
|
|
|
/* BINOP = / */
|
|
lgot = 1LL << 32;
|
|
lexp = 1LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot /= expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL << 32;
|
|
lexp = 1LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot /= expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL << 32;
|
|
lexp = 1LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = lgot / expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 2LL;
|
|
lexp = 2LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL << N;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = expr / lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 2LL;
|
|
lexp = 2LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL << N;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = lgot / expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 2LL;
|
|
lexp = 2LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL << N;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = expr / lgot; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
/* BINOP = & */
|
|
lgot = ~0LL;
|
|
lexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot &= expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = ~0LL;
|
|
iexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot &= expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = ~0LL;
|
|
lexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = lgot & expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = ~0LL;
|
|
lexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = expr & lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = ~0LL;
|
|
iexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = lgot & expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = ~0LL;
|
|
lexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = expr & lgot; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
/* BINOP = ^ */
|
|
lgot = ~0LL;
|
|
lexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1 << i;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot ^= expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = ~0LL;
|
|
iexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot ^= expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = ~0LL;
|
|
lexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = lgot ^ expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = ~0LL;
|
|
lexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = expr ^ lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = ~0LL;
|
|
iexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = lgot ^ expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = ~0LL;
|
|
lexp = 0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = expr ^ lgot; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
/* BINOP = | */
|
|
lgot = 0LL;
|
|
lexp = ~0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1 << i;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot |= expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 0LL;
|
|
iexp = ~0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot |= expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 0LL;
|
|
lexp = ~0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = lgot | expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 0LL;
|
|
lexp = ~0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = expr | lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 0LL;
|
|
iexp = ~0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = lgot | expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 0LL;
|
|
lexp = ~0LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = ~(1 << i);
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = expr | lgot; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
/* BINOP = << */
|
|
lgot = 1LL;
|
|
lexp = 1LL << N;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot <<= expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
iexp = 1LL << N;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot <<= expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 1LL << N;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = lgot << expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 2LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < 1; i++)
|
|
{
|
|
long long expr = 1LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = expr << lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 2LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < 1; i++)
|
|
{
|
|
long long expr = 1LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = lgot << expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 2LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < 1; i++)
|
|
{
|
|
long long expr = 1LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = expr << lgot; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
/* BINOP = >> */
|
|
lgot = 1LL << N;
|
|
lexp = 1LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot >>= expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL << N;
|
|
iexp = 1LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot >>= expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL << N;
|
|
lexp = 1LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = lgot >> expr; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 1LL << (N - 1);
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < 1; i++)
|
|
{
|
|
long long expr = 1LL << N;
|
|
|
|
#pragma acc atomic capture
|
|
{ ldata[i] = lgot; lgot = expr >> lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL << N;
|
|
lexp = 1LL;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = lgot >> expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 1LL << (N - 1);
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < 1; i++)
|
|
{
|
|
long long expr = 1LL << N;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = expr >> lgot; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
// FLOAT FLOAT FLOAT
|
|
|
|
/* BINOP = + */
|
|
fgot = 0.0;
|
|
fexp = 32.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot += expr; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 0.0;
|
|
fexp = 32.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot += expr; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 0.0;
|
|
fexp = 32.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ idata[i] = fgot; fgot = fgot + expr; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 0.0;
|
|
fexp = 32.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot = expr + fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 0.0;
|
|
fexp = 32.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot = fgot + expr; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 0.0;
|
|
fexp = 32.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot = expr + fgot; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
/* BINOP = * */
|
|
fgot = 1.0;
|
|
fexp = 8192.0*8192.0*64.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 2.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot *= expr; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 1.0;
|
|
fexp = 8192.0*8192.0*64.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 2.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot *= expr; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 1.0;
|
|
fexp = 8192.0*8192.0*64.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 2.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot = fgot * expr; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 1.0;
|
|
fexp = 8192.0*8192.0*64.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 2.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot = expr * fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
lgot = 1LL;
|
|
lexp = 1LL << 32;
|
|
|
|
#pragma acc data copy (lgot, ldata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2LL;
|
|
|
|
#pragma acc atomic capture
|
|
{ lgot = lgot * expr; ldata[i] = lgot; }
|
|
}
|
|
}
|
|
|
|
if (lexp != lgot)
|
|
abort ();
|
|
|
|
fgot = 1.0;
|
|
fexp = 8192.0*8192.0*64.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 2;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot = expr * fgot; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
/* BINOP = - */
|
|
fgot = 32.0;
|
|
fexp = 0.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot -= expr; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 32.0;
|
|
fexp = 0.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot -= expr; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 32.0;
|
|
fexp = 0.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot = fgot - expr; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 1.0;
|
|
fexp = 1.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot = expr - fgot; }
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
if (i % 2 == 0)
|
|
{
|
|
if (fdata[i] != 1.0)
|
|
abort ();
|
|
}
|
|
else
|
|
{
|
|
if (fdata[i] != 0.0)
|
|
abort ();
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 1.0;
|
|
fexp = -31.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot = fgot - expr; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 1.0;
|
|
fexp = 1.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot = expr - fgot; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < N; i++)
|
|
if (i % 2 == 0)
|
|
{
|
|
if (fdata[i] != 0.0)
|
|
abort ();
|
|
}
|
|
else
|
|
{
|
|
if (fdata[i] != 1.0)
|
|
abort ();
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
/* BINOP = / */
|
|
fgot = 8192.0*8192.0*64.0;
|
|
fexp = 1.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 2.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot /= expr; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 8192.0*8192.0*64.0;
|
|
fexp = 1.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 2.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot /= expr; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 8192.0*8192.0*64.0;
|
|
fexp = 1.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 2.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot = fgot / expr; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 8192.0*8192.0*64.0;
|
|
fexp = 1.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 1.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fdata[i] = fgot; fgot = expr / fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 4.0;
|
|
fexp = 4.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
long long expr = 1LL << N;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot = fgot / expr; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
fgot = 4.0;
|
|
fexp = 4.0;
|
|
|
|
#pragma acc data copy (fgot, fdata[0:N])
|
|
{
|
|
#pragma acc parallel loop
|
|
for (i = 0; i < N; i++)
|
|
{
|
|
float expr = 2.0;
|
|
|
|
#pragma acc atomic capture
|
|
{ fgot = expr / fgot; fdata[i] = fgot; }
|
|
}
|
|
}
|
|
|
|
if (fexp != fgot)
|
|
abort ();
|
|
|
|
return 0;
|
|
}
|