OpenE2K/gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity
6e0df3b7a8
gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c
Thomas Schwinge 4bf9e5a8a2 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

867 lines
13 KiB
C
Raw Blame History

/* { dg-do run } */
#include <stdlib.h>
int
main(int argc, char **argv)
{
int iexp, igot;
long long lexp, lgot;
int N = 32;
int idata[N];
long long ldata[N];
float fexp, fgot;
float fdata[N];
int i;
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++;
}
}
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--;
}
}
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;
}
}
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;
}
}
/* 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 += 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;
}
}
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 = expr + igot;
}
}
if (iexp != igot)
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 = 2LL;
#pragma acc atomic capture
ldata[i] = lgot *= expr;
}
}
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 = 2LL;
#pragma acc atomic capture
ldata[i] = lgot = lgot * expr;
}
}
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 = 2LL;
#pragma acc atomic capture
ldata[i] = lgot = expr * 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 -= 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
idata[i] = igot = igot - expr;
}
}
if (iexp != igot)
abort ();
igot = 32;
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 = expr - igot;
}
}
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 /= 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
ldata[i] = 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 = expr / lgot;
}
}
if (lexp != lgot)
abort ();
/* BINOP = & */
igot = ~0;
iexp = 0;
#pragma acc data copy (igot, idata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = 1 << i;
#pragma acc atomic capture
idata[i] = igot &= expr;
}
}
if (iexp != igot)
abort ();
igot = ~0;
iexp = 0;
#pragma acc data copy (igot, idata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = 1 << i;
#pragma acc atomic capture
idata[i] = igot = igot & expr;
}
}
if (iexp != igot)
abort ();
igot = ~0;
iexp = 0;
#pragma acc data copy (igot, idata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = 1 << i;
#pragma acc atomic capture
idata[i] = igot = expr & igot;
}
}
if (iexp != igot)
abort ();
/* BINOP = ^ */
igot = ~0;
iexp = 0;
#pragma acc data copy (igot, idata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = 1 << i;
#pragma acc atomic capture
idata[i] = igot ^= expr;
}
}
if (iexp != igot)
abort ();
igot = ~0;
iexp = 0;
#pragma acc data copy (igot, idata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = 1 << i;
#pragma acc atomic capture
idata[i] = igot = igot ^ expr;
}
}
if (iexp != igot)
abort ();
igot = ~0;
iexp = 0;
#pragma acc data copy (igot, idata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = 1 << i;
#pragma acc atomic capture
idata[i] = igot = expr ^ igot;
}
}
if (iexp != igot)
abort ();
/* BINOP = | */
igot = 0;
iexp = ~0;
#pragma acc data copy (igot, idata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = 1 << i;
#pragma acc atomic capture
idata[i] = igot |= expr;
}
}
if (iexp != igot)
abort ();
igot = 0;
iexp = ~0;
#pragma acc data copy (igot, idata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = 1 << i;
#pragma acc atomic capture
idata[i] = igot = igot | expr;
}
}
if (iexp != igot)
abort ();
igot = 0;
iexp = ~0;
#pragma acc data copy (igot, idata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
int expr = 1 << i;
#pragma acc atomic capture
idata[i] = igot = expr | igot;
}
}
if (iexp != igot)
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++)
{
int expr = 1;
#pragma acc atomic capture
ldata[i] = lgot <<= expr;
}
}
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++)
{
int expr = 1;
#pragma acc atomic capture
idata[i] = lgot = lgot << expr;
}
}
if (lexp != lgot)
abort ();
lgot = 1LL;
lexp = 2LL;
#pragma acc data copy (lgot, ldata[0:N])
{
#pragma acc parallel
{
long long expr = 1LL;
#pragma acc atomic capture
ldata[0] = lgot = expr << 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 >>= expr;
}
}
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++)
{
int expr = 1;
#pragma acc atomic capture
ldata[i] = lgot = lgot >> expr;
}
}
if (lexp != lgot)
abort ();
lgot = 1LL << 63;
lexp = 1LL << 32;
#pragma acc data copy (lgot, ldata[0:N])
{
#pragma acc parallel
{
long long expr = 1LL << 32;
#pragma acc atomic capture
ldata[0] = lgot = expr >> lgot;
}
}
if (lexp != lgot)
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++)
{
#pragma acc atomic capture
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++)
{
#pragma acc atomic capture
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++)
{
#pragma acc atomic capture
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++)
{
#pragma acc atomic capture
fdata[i] = --fgot;
}
}
if (fexp != fgot)
abort ();
/* 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 += 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;
}
}
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 = expr + 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 *= 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++)
{
long long expr = 2LL;
#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
fdata[i] = fgot = expr * 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 -= 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
fdata[i] = fgot = fgot - expr;
}
}
if (fexp != fgot)
abort ();
fgot = 1.0;
fexp = 0.0;
#pragma acc data copy (fgot, fdata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
float expr = 32.0;
#pragma acc atomic capture
fdata[i] = fgot = expr - fgot;
}
}
for (i = 0; i < N; i++)
if (i % 2 == 0)
{
if (fdata[i] != 31.0)
abort ();
}
else
{
if (fdata[i] != 1.0)
abort ();
}
/* BINOP = / */
fexp = 1.0;
fgot = 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 /= expr;
}
}
if (fexp != fgot)
abort ();
fexp = 1.0;
fgot = 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 ();
fexp = 1.0;
fgot = 8192.0*8192.0*64.0;
#pragma acc data copy (fgot, fdata[0:N])
{
#pragma acc parallel
{
float expr = 8192.0*8192.0*64.0;
#pragma acc atomic capture
fdata[0] = fgot = expr / fgot;
}
}
if (fexp != fgot)
abort ();
return 0;
}
Reference in New Issue View Git Blame Copy Permalink