FreeChainXenon/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions
gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_update-1.c

761 lines
11 KiB
C
Raw Normal View History

OpenACC atomic directive 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
2015-11-03 12:28:22 +01:00
/* { dg-do run } */
#include <stdlib.h>
int
main(int argc, char **argv)
{
float fexp, fgot;
int iexp, igot;
long long lexp, lgot;
int N = 32;
int i;
fgot = 1234.0;
fexp = 1235.0;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < 1; i++)
#pragma acc atomic update
fgot++;
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = fgot - N;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
#pragma acc atomic update
fgot--;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = fgot + N;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
#pragma acc atomic update
++fgot;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = fgot - N;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
#pragma acc atomic update
--fgot;
}
}
if (fexp != fgot)
abort ();
/* BINOP = + */
fgot = 1234.0;
fexp = fgot + N;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 1.0;
#pragma acc atomic update
fgot += expr;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = fgot + N;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 1.0;
#pragma acc atomic update
fgot = fgot + expr;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = fgot + N;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 1.0;
#pragma acc atomic update
fgot = expr + fgot;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 0.5;
#pragma acc atomic update
fgot = (expr + expr) + fgot;
}
}
if (fexp != fgot)
abort ();
/* BINOP = * */
fgot = 1234.0;
fexp = 1234.0;
for (i = 0; i < N; i++)
fexp *= 2.0;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 2.0;
#pragma acc atomic update
fgot *= expr;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = 1234.0;
for (i = 0; i < N; i++)
fexp = fexp * 2.0;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 2.0;
#pragma acc atomic update
fgot = fgot * expr;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = 1234.0;
for (i = 0; i < N; i++)
fexp = 2.0 * fexp;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 2.0;
#pragma acc atomic update
fgot = expr * fgot;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 1.0;
#pragma acc atomic update
fgot = (expr + expr) * fgot;
}
}
if (fexp != fgot)
abort ();
/* BINOP = - */
fgot = 1234.0;
fexp = 1234.0;
for (i = 0; i < N; i++)
fexp -= 2.0;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 2.0;
#pragma acc atomic update
fgot -= expr;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = 1234.0;
for (i = 0; i < N; i++)
fexp = fexp - 2.0;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 2.0;
#pragma acc atomic update
fgot = fgot - expr;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = 1234.0;
for (i = 0; i < N; i++)
fexp = 2.0 - fexp;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 2.0;
#pragma acc atomic update
fgot = expr - fgot;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 1.0;
#pragma acc atomic update
fgot = (expr + expr) - fgot;
}
}
if (fexp != fgot)
abort ();
/* BINOP = / */
fgot = 1234.0;
fexp = 1234.0;
for (i = 0; i < N; i++)
fexp /= 2.0;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 2.0;
#pragma acc atomic update
fgot /= expr;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = 1234.0;
for (i = 0; i < N; i++)
fexp = fexp / 2.0;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 2.0;
#pragma acc atomic update
fgot = fgot / expr;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = 1234.0;
for (i = 0; i < N; i++)
fexp = 2.0 / fexp;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 2.0;
#pragma acc atomic update
fgot = expr / fgot;
}
}
if (fexp != fgot)
abort ();
fgot = 1234.0;
fexp = 1234.0;
for (i = 0; i < N; i++)
fexp = 2.0 / fexp;
#pragma acc data copy (fgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 1.0;
#pragma acc atomic update
fgot = (expr + expr) / fgot;
}
}
if (fexp != fgot)
abort ();
/* BINOP = & */
igot = ~0;
iexp = 0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = ~(1 << i);
#pragma acc atomic update
igot &= expr;
}
}
if (iexp != igot)
abort ();
igot = ~0;
iexp = 0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = ~(1 << i);
#pragma acc atomic update
igot = igot / expr;
}
}
if (iexp != igot)
abort ();
igot = ~0;
iexp = 0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = ~(1 << i);
#pragma acc atomic update
igot = expr & igot;
}
}
if (iexp != igot)
abort ();
igot = ~0;
iexp = 0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = ~(1 << i);
int zero = 0;
#pragma acc atomic update
igot = (expr + zero) & igot;
}
}
if (iexp != igot)
abort ();
/* BINOP = ^ */
igot = ~0;
iexp = 0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = (1 << i);
#pragma acc atomic update
igot ^= expr;
}
}
if (iexp != igot)
abort ();
igot = ~0;
iexp = 0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = (1 << i);
#pragma acc atomic update
igot = igot ^ expr;
}
}
if (iexp != igot)
abort ();
igot = ~0;
iexp = 0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = (1 << i);
#pragma acc atomic update
igot = expr ^ igot;
}
}
if (iexp != igot)
abort ();
igot = ~0;
iexp = 0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = (1 << i);
int zero = 0;
#pragma acc atomic update
igot = (expr + zero) ^ igot;
}
}
if (iexp != igot)
abort ();
/* BINOP = | */
igot = 0;
iexp = ~0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = (1 << i);
#pragma acc atomic update
igot |= expr;
}
}
if (iexp != igot)
abort ();
igot = 0;
iexp = ~0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = (1 << i);
#pragma acc atomic update
igot = igot | expr;
}
}
if (iexp != igot)
abort ();
igot = 0;
iexp = ~0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = (1 << i);
#pragma acc atomic update
igot = expr | igot;
}
}
if (iexp != igot)
abort ();
igot = 0;
iexp = ~0;
#pragma acc data copy (igot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = (1 << i);
int zero = 0;
#pragma acc atomic update
igot = (expr + zero) | igot;
}
}
if (iexp != igot)
abort ();
/* BINOP = << */
lgot = 1LL;
lexp = 1LL << N;
#pragma acc data copy (lgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
long long expr = 1LL;
#pragma acc atomic update
lgot <<= expr;
}
}
if (lexp != lgot)
abort ();
lgot = 1LL;
lexp = 1LL << N;
#pragma acc data copy (lgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
long long expr = 1LL;
#pragma acc atomic update
lgot = lgot << expr;
}
}
if (lexp != lgot)
abort ();
lgot = 1LL;
lexp = 2LL;
#pragma acc data copy (lgot)
{
#pragma acc parallel loop
for (i = 0; i < 1; i++)
{
long long expr = 1LL;
#pragma acc atomic update
lgot = expr << lgot;
}
}
if (lexp != lgot)
abort ();
lgot = 1LL;
lexp = 2LL;
#pragma acc data copy (lgot)
{
#pragma acc parallel loop
for (i = 0; i < 1; i++)
{
long long expr = 1LL;
long long zero = 0LL;
#pragma acc atomic update
lgot = (expr + zero) << lgot;
}
}
if (lexp != lgot)
abort ();
/* BINOP = >> */
lgot = 1LL << N;
lexp = 1LL;
#pragma acc data copy (lgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
long long expr = 1LL;
#pragma acc atomic update
lgot >>= expr;
}
}
if (lexp != lgot)
abort ();
lgot = 1LL << N;
lexp = 1LL;
#pragma acc data copy (lgot)
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
long long expr = 1LL;
#pragma acc atomic update
lgot = lgot >> expr;
}
}
if (lexp != lgot)
abort ();
lgot = 1LL;
lexp = 1LL << (N - 1);
#pragma acc data copy (lgot)
{
#pragma acc parallel loop
for (i = 0; i < 1; i++)
{
long long expr = 1LL << N;
#pragma acc atomic update
lgot = expr >> lgot;
}
}
if (lexp != lgot)
abort ();
lgot = 1LL;
lexp = 1LL << (N - 1);
#pragma acc data copy (lgot)
{
#pragma acc parallel loop
for (i = 0; i < 1; i++)
{
long long expr = 1LL << N;
long long zero = 0LL;
#pragma acc atomic update
lgot = (expr + zero) >> lgot;
}
}
if (lexp != lgot)
abort ();
return 0;
}
Reference in a new issue Copy permalink