gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/data-2.c
Cesar Philippidis e46c777050 c-parser.c (c_parser_oacc_declare): Add support for GOMP_MAP_FIRSTPRIVATE_POINTER.
gcc/c/
	* c-parser.c (c_parser_oacc_declare): Add support for
	GOMP_MAP_FIRSTPRIVATE_POINTER.
	* c-typeck.c (handle_omp_array_sections_1): Replace bool is_omp
	argument with enum c_omp_region_type ort.
	(handle_omp_array_sections): Likewise.  Update call to
	handle_omp_array_sections_1.
	(c_finish_omp_clauses): Add specific errors and warning messages for
	OpenACC.  Use firsrtprivate pointers for OpenACC subarrays.  Update
	call to handle_omp_array_sections.

	gcc/cp/
	* parser.c (cp_parser_oacc_declare): Add support for
	GOMP_MAP_FIRSTPRIVATE_POINTER.
	* semantics.c (handle_omp_array_sections_1): Replace bool is_omp
	argument with enum c_omp_region_type ort.  Don't privatize OpenACC
	non-static members.
	(handle_omp_array_sections): Replace bool is_omp argument with enum
	c_omp_region_type ort.  Update call to handle_omp_array_sections_1.
	(finish_omp_clauses): Add specific errors and warning messages for
	OpenACC.  Use firsrtprivate pointers for OpenACC subarrays.  Update
	call to handle_omp_array_sections.

	gcc/
	* gimplify.c (omp_notice_variable): Use zero-length arrays for data
	pointers inside OACC_DATA regions.
	(gimplify_scan_omp_clauses): Prune firstprivate clause associated
	with OACC_DATA, OACC_ENTER_DATA and OACC_EXIT data regions.
	(gimplify_adjust_omp_clauses): Fix typo in comment.

	gcc/testsuite/
	* c-c++-common/goacc/data-clause-duplicate-1.c: Adjust test.
	* c-c++-common/goacc/deviceptr-1.c: Likewise.
	* c-c++-common/goacc/kernels-alias-3.c: Likewise.
	* c-c++-common/goacc/kernels-alias-4.c: Likewise.
	* c-c++-common/goacc/kernels-alias-5.c: Likewise.
	* c-c++-common/goacc/kernels-alias-8.c: Likewise.
	* c-c++-common/goacc/kernels-alias-ipa-pta-3.c: Likewise.
	* c-c++-common/goacc/pcopy.c: Likewise.
	* c-c++-common/goacc/pcopyin.c: Likewise.
	* c-c++-common/goacc/pcopyout.c: Likewise.
	* c-c++-common/goacc/pcreate.c: Likewise.
	* c-c++-common/goacc/pr70688.c: New test.
	* c-c++-common/goacc/present-1.c: Adjust test.
	* c-c++-common/goacc/reduction-5.c: Likewise.
	* g++.dg/goacc/data-1.C: New test.

	libgomp/
	* oacc-mem.c (acc_malloc): Update handling of shared-memory targets.
	(acc_free): Likewise.
	(acc_memcpy_to_device): Likewise.
	(acc_memcpy_from_device): Likewise.
	(acc_deviceptr): Likewise.
	(acc_hostptr): Likewise.
	(acc_is_present): Likewise.
	(acc_map_data): Likewise.
	(acc_unmap_data): Likewise.
	(present_create_copy): Likewise.
	(delete_copyout): Likewise.
	(update_dev_host): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Remove xfail.
	* testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/data-2.c: Adjust test.
	* testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/enter_exit-lib.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/lib-13.c: Adjust test so that
	it only runs on nvptx targets.
	* testsuite/libgomp.oacc-c-c++-common/lib-14.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-15.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-17.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-18.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-20.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-21.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-22.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-23.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-24.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-28.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-29.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-30.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-34.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-42.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-43.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-44.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-47.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-48.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-52.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-53.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/lib-54.c: Likewise.

From-SVN: r236678
2016-05-24 15:54:21 -07:00

166 lines
3.5 KiB
C

/* Test 'acc enter/exit data' regions. */
/* { dg-do run } */
#include <stdlib.h>
int
main (int argc, char **argv)
{
int N = 128; //1024 * 1024;
float *a, *b, *c, *d, *e;
int i;
int nbytes;
nbytes = N * sizeof (float);
a = (float *) malloc (nbytes);
b = (float *) malloc (nbytes);
c = (float *) malloc (nbytes);
d = (float *) malloc (nbytes);
e = (float *) malloc (nbytes);
for (i = 0; i < N; i++)
{
a[i] = 3.0;
b[i] = 0.0;
}
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async
#pragma acc parallel present (a[0:N], b[0:N]) async wait
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = a[i];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async
#pragma acc wait
for (i = 0; i < N; i++)
{
if (a[i] != 3.0)
abort ();
if (b[i] != 3.0)
abort ();
}
for (i = 0; i < N; i++)
{
a[i] = 2.0;
b[i] = 0.0;
}
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (N) async (1)
#pragma acc parallel present (a[0:N], b[0:N]) async (1)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = a[i];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait (1) async (1)
#pragma acc wait (1)
for (i = 0; i < N; i++)
{
if (a[i] != 2.0)
abort ();
if (b[i] != 2.0)
abort ();
}
for (i = 0; i < N; i++)
{
a[i] = 3.0;
b[i] = 0.0;
c[i] = 0.0;
d[i] = 0.0;
}
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (N) async (1)
#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1)
#pragma acc loop
for (i = 0; i < N; i++)
b[i] = (a[i] * a[i] * a[i]) / a[i];
#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1)
#pragma acc loop
for (i = 0; i < N; i++)
c[i] = (a[i] + a[i] + a[i] + a[i]) / a[i];
#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1)
#pragma acc loop
for (i = 0; i < N; i++)
d[i] = ((a[i] * a[i] + a[i]) / a[i]) - a[i];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) copyout (d[0:N]) wait (1, 2, 3) async (1)
#pragma acc wait (1)
for (i = 0; i < N; i++)
{
if (a[i] != 3.0)
abort ();
if (b[i] != 9.0)
abort ();
if (c[i] != 4.0)
abort ();
if (d[i] != 1.0)
abort ();
}
for (i = 0; i < N; i++)
{
a[i] = 2.0;
b[i] = 0.0;
c[i] = 0.0;
d[i] = 0.0;
e[i] = 0.0;
}
#pragma acc enter data copyin (a[0:N]) copyin (b[0:N]) copyin (c[0:N]) copyin (d[0:N]) copyin (e[0:N]) copyin (N) async (1)
#pragma acc parallel present (a[0:N], b[0:N]) async (1) wait (1)
for (int ii = 0; ii < N; ii++)
b[ii] = (a[ii] * a[ii] * a[ii]) / a[ii];
#pragma acc parallel present (a[0:N], c[0:N]) async (2) wait (1)
for (int ii = 0; ii < N; ii++)
c[ii] = (a[ii] + a[ii] + a[ii] + a[ii]) / a[ii];
#pragma acc parallel present (a[0:N], d[0:N]) async (3) wait (1)
for (int ii = 0; ii < N; ii++)
d[ii] = ((a[ii] * a[ii] + a[ii]) / a[ii]) - a[ii];
#pragma acc parallel present (a[0:N], b[0:N], c[0:N], d[0:N], e[0:N]) \
wait (1) async (4)
for (int ii = 0; ii < N; ii++)
e[ii] = a[ii] + b[ii] + c[ii] + d[ii];
#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) copyout (c[0:N]) \
copyout (d[0:N]) copyout (e[0:N]) wait (1, 2, 3, 4) async (1)
#pragma acc wait (1)
for (i = 0; i < N; i++)
{
if (a[i] != 2.0)
abort ();
if (b[i] != 4.0)
abort ();
if (c[i] != 4.0)
abort ();
if (d[i] != 1.0)
abort ();
if (e[i] != 11.0)
abort ();
}
return 0;
}