gcc/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c
Julian Brown 6f5b4b64d2 openacc: Adjust dynamic reference count semantics
This patch adjusts how dynamic reference counts work so that they match
the semantics of the source program more closely, instead of representing
"excess" reference counts beyond those that represent pointers in the
internal libgomp splay-tree data structure. This allows some corner
cases to be handled more gracefully.

2020-07-10  Julian Brown  <julian@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

	libgomp/
	* libgomp.h (struct splay_tree_key_s): Change virtual_refcount to
	dynamic_refcount.
	(struct gomp_device_descr): Remove GOMP_MAP_VARS_OPENACC_ENTER_DATA.
	* oacc-mem.c (acc_map_data): Substitute virtual_refcount for
	dynamic_refcount.
	(acc_unmap_data): Update comment.
	(goacc_map_var_existing, goacc_enter_datum): Adjust for
	dynamic_refcount semantics.
	(goacc_exit_datum_1, goacc_exit_datum): Re-add some error checking.
	Adjust for dynamic_refcount semantics.
	(goacc_enter_data_internal): Implement "present" case of dynamic
	memory-map handling here.  Update "non-present" case for
	dynamic_refcount semantics.
	(goacc_exit_data_internal): Use goacc_exit_datum_1.
	* target.c (gomp_map_vars_internal): Remove
	GOMP_MAP_VARS_OPENACC_ENTER_DATA handling.  Update for dynamic_refcount
	handling.
	(gomp_unmap_vars_internal): Remove virtual_refcount handling.
	(gomp_load_image_to_device): Substitute dynamic_refcount for
	virtual_refcount.
	* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: Remove XFAILs.
	* testsuite/libgomp.oacc-c-c++-common/refcounting-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/refcounting-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/struct-3-1-1.c: New test.
	* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: Remove XFAILs and
	trace output.
	* testsuite/libgomp.oacc-fortran/deep-copy-6-no_finalize.F90: Remove
	trace output.
	* testsuite/libgomp.oacc-fortran/dynamic-incr-structural-1.f90: New
	test.
	* testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c:
	Remove stale comment.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-1.f90: Remove XFAILs.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-1-2.F90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-2-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-3-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/mdc-refcount-1-4-1.f90: Adjust XFAIL.

Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
2020-07-10 08:07:12 -07:00

179 lines
3.9 KiB
C

/* Verify that 'acc_copyout' etc. is a no-op if there's still a structured
reference count. */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <assert.h>
#include <stdlib.h>
#include <openacc.h>
const int c0 = 58;
const int c1 = 81;
static void
assign_array (char *array, size_t size, char value)
{
for (size_t i = 0; i < size; ++i)
array[i] = value;
}
static void
verify_array (const char *array, size_t size, char value)
{
for (size_t i = 0; i < size; ++i)
assert (array[i] == value);
}
float global_var;
#pragma acc declare create (global_var)
static void
test_acc_declare ()
{
assert (acc_is_present (&global_var, sizeof global_var));
global_var = c0;
#pragma acc update device (global_var)
global_var = c1;
acc_copyout (&global_var, sizeof global_var);
assert (acc_is_present (&global_var, sizeof global_var));
assert (global_var == c1);
global_var = c1;
acc_copyout_finalize (&global_var, sizeof global_var);
assert (acc_is_present (&global_var, sizeof global_var));
assert (global_var == c1);
void *global_var_d_p = acc_deviceptr (&global_var);
assert (global_var_d_p);
void *d_p = acc_copyin (&global_var, sizeof global_var);
assert (d_p == global_var_d_p);
acc_copyout (&global_var, sizeof global_var);
assert (acc_is_present (&global_var, sizeof global_var));
d_p = acc_copyin (&global_var, sizeof global_var);
assert (d_p == global_var_d_p);
d_p = acc_copyin (&global_var, sizeof global_var);
assert (d_p == global_var_d_p);
global_var = c1;
acc_copyout_finalize (&global_var, sizeof global_var);
assert (acc_is_present (&global_var, sizeof global_var));
assert (global_var == c1);
global_var = c1;
acc_copyout (&global_var, sizeof global_var);
assert (acc_is_present (&global_var, sizeof global_var));
assert (global_var == c1);
}
static void
test_acc_map_data ()
{
const int N = 801;
char *h = (char *) malloc (N);
assert (h);
void *d = acc_malloc (N);
assert (d);
acc_map_data (h, d, N);
assert (acc_is_present (h, N));
assign_array (h, N, c0);
#pragma acc update device (h[0:N])
assign_array (h, N, c1);
#pragma acc exit data copyout (h[0:N])
assert (acc_is_present (h, N));
verify_array (h, N, c1);
assign_array (h, N, c1);
#pragma acc exit data copyout (h[0:N]) finalize
assert (acc_is_present (h, N));
verify_array (h, N, c1);
#pragma acc enter data copyin (h[0:N])
assign_array (h, N, c1);
#pragma acc exit data copyout (h[0:N])
assert (acc_is_present (h, N));
verify_array (h, N, c1);
#pragma acc enter data copyin (h[0:N])
#pragma acc enter data copyin (h[0:N])
assign_array (h, N, c1);
#pragma acc exit data copyout (h[0:N]) finalize
assert (acc_is_present (h, N));
verify_array (h, N, c1);
assign_array (h, N, c1);
#pragma acc exit data copyout (h[0:N])
assert (acc_is_present (h, N));
verify_array (h, N, c1);
}
static void
test_acc_data ()
{
#define N 23
char h[N];
assign_array (h, N, c0);
#pragma acc data copyin (h)
{
assert (acc_is_present (h, sizeof h));
assign_array (h, N, c1);
acc_copyout_finalize (h, sizeof h);
assert (acc_is_present (h, sizeof h));
verify_array (h, N, c1);
assign_array (h, N, c1);
acc_copyout (h, sizeof h);
assert (acc_is_present (h, sizeof h));
verify_array (h, N, c1);
acc_copyin (h, sizeof h);
assign_array (h, N, c1);
acc_copyout (h, sizeof h);
assert (acc_is_present (h, sizeof h));
verify_array (h, N, c1);
acc_copyin (h, sizeof h);
acc_copyin (h, sizeof h);
assign_array (h, N, c1);
acc_copyout_finalize (h, sizeof h);
assert (acc_is_present (h, sizeof h));
verify_array (h, N, c1);
assign_array (h, N, c1);
acc_copyout (h, sizeof h);
assert (acc_is_present (h, sizeof h));
verify_array (h, N, c1);
}
#undef N
}
int
main ()
{
test_acc_declare ();
test_acc_map_data ();
test_acc_data ();
return 0;
}