OpenE2K/gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[PR92843] [OpenACC] Fix dynamic reference counting for structured 'REFCOUNT_INFINITY'

Browse Source 
libgomp/
	PR libgomp/92843
	* oacc-mem.c (present_create_copy, delete_copyout): Fix dynamic
	reference counting for structured 'REFCOUNT_INFINITY'.  Add some
	assertions.
	(goacc_insert_pointer, goacc_remove_pointer): Adjust accordingly.
	* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: New file.
	* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Fix OpenACC.
	* testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.

From-SVN: r279234
This commit is contained in:
Thomas Schwinge 2019-12-11 17:49:27 +01:00 committed by Thomas Schwinge
parent 57963e3934
commit d6e8c01cff
6 changed files with 242 additions and 21 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
libgomp/ChangeLog
View File

@ -1,5 +1,15 @@
2019-12-11 Thomas Schwinge <thomas@codesourcery.com>
PR libgomp/92843
* oacc-mem.c (present_create_copy, delete_copyout): Fix dynamic
reference counting for structured 'REFCOUNT_INFINITY'. Add some
assertions.
(goacc_insert_pointer, goacc_remove_pointer): Adjust accordingly.
* testsuite/libgomp.oacc-c-c++-common/pr92843-1.c: New file.
* testsuite/libgomp.oacc-c-c++-common/clauses-1.c: Fix OpenACC.
* testsuite/libgomp.oacc-c-c++-common/lib-82.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/nested-1.c: Likewise.
* oacc-parallel.c (find_pointer, GOACC_enter_exit_data): Move...
* oacc-mem.c: ... here.
(gomp_acc_insert_pointer, gomp_acc_remove_pointer): Rename to

42
libgomp/oacc-mem.c
View File

@ -543,11 +543,11 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s);
}
assert (n->refcount != REFCOUNT_LINK);
if (n->refcount != REFCOUNT_INFINITY)
{
n->refcount++;
n->dynamic_refcount++;
}
n->refcount++;
n->dynamic_refcount++;
gomp_mutex_unlock (&acc_dev->lock);
}
else if (!(f & FLAG_CREATE))
@ -573,8 +573,10 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
tgt = gomp_map_vars_async (acc_dev, aq, mapnum, &hostaddrs, NULL, &s,
&kinds, true, GOMP_MAP_VARS_OPENACC);
/* Initialize dynamic refcount. */
tgt->list[0].key->dynamic_refcount = 1;
n = tgt->list[0].key;
assert (n->refcount == 1);
assert (n->dynamic_refcount == 0);
n->dynamic_refcount++;
d = tgt->to_free;
}
@ -698,12 +700,9 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
(void *) h, (int) s, (void *) n->host_start, (int) host_size);
}
if (n->refcount == REFCOUNT_INFINITY)
{
n->refcount = 0;
n->dynamic_refcount = 0;
}
if (n->refcount < n->dynamic_refcount)
assert (n->refcount != REFCOUNT_LINK);
if (n->refcount != REFCOUNT_INFINITY
&& n->refcount < n->dynamic_refcount)
{
gomp_mutex_unlock (&acc_dev->lock);
gomp_fatal ("Dynamic reference counting assert fail\n");
@ -711,13 +710,15 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
if (f & FLAG_FINALIZE)
{
n->refcount -= n->dynamic_refcount;
if (n->refcount != REFCOUNT_INFINITY)
n->refcount -= n->dynamic_refcount;
n->dynamic_refcount = 0;
}
else if (n->dynamic_refcount)
{
if (n->refcount != REFCOUNT_INFINITY)
n->refcount--;
n->dynamic_refcount--;
n->refcount--;
}
if (n->refcount == 0)
@ -895,6 +896,8 @@ goacc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
splay_tree_key n;
gomp_mutex_lock (&acc_dev->lock);
n = lookup_host (acc_dev, *hostaddrs, *sizes);
assert (n->refcount != REFCOUNT_INFINITY
&& n->refcount != REFCOUNT_LINK);
gomp_mutex_unlock (&acc_dev->lock);
tgt = n->tgt;
@ -917,10 +920,11 @@ goacc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
goacc_aq aq = get_goacc_asyncqueue (async);
tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs,
NULL, sizes, kinds, true, GOMP_MAP_VARS_OPENACC);
splay_tree_key n = tgt->list[0].key;
assert (n->refcount == 1);
assert (n->dynamic_refcount == 0);
n->dynamic_refcount++;
gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
/* Initialize dynamic refcount. */
tgt->list[0].key->dynamic_refcount = 1;
}
static void
@ -950,6 +954,8 @@ goacc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
t = n->tgt;
assert (n->refcount != REFCOUNT_INFINITY
&& n->refcount != REFCOUNT_LINK);
if (n->refcount < n->dynamic_refcount)
{
gomp_mutex_unlock (&acc_dev->lock);
@ -963,8 +969,8 @@ goacc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
}
else if (n->dynamic_refcount)
{
n->dynamic_refcount--;
n->refcount--;
n->dynamic_refcount--;
}
gomp_mutex_unlock (&acc_dev->lock);

16
libgomp/testsuite/libgomp.oacc-c-c++-common/clauses-1.c
View File

@ -469,7 +469,9 @@ main (int argc, char **argv)
if (!acc_is_present (c, (N * sizeof (float))))
abort ();
acc_copyout (b, N * sizeof (float));
d = (float *) acc_deviceptr (b);
acc_memcpy_from_device (b, d, N * sizeof (float));
for (i = 0; i < N; i++)
{
@ -485,10 +487,22 @@ main (int argc, char **argv)
if (acc_is_present (a, N * sizeof (float)))
abort ();
d = (float *) acc_deviceptr (b);
acc_unmap_data (b);
if (acc_is_present (b, N * sizeof (float)))
abort ();
acc_free (d);
d = (float *) acc_deviceptr (c);
acc_unmap_data (c);
if (acc_is_present (c, N * sizeof (float)))
abort ();
acc_free (d);
for (i = 0; i < N; i++)

6
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-82.c
View File

@ -120,9 +120,13 @@ main (int argc, char **argv)
for (i = 0; i < N; i++)
{
acc_copyout (a[i], nbytes);
acc_memcpy_from_device (a[i], d_a[i], nbytes);
if (*a[i] != i)
abort ();
acc_unmap_data (a[i]);
acc_free (d_a[i]);
}
free (streams);

10
libgomp/testsuite/libgomp.oacc-c-c++-common/nested-1.c
View File

@ -517,7 +517,9 @@ main (int argc, char **argv)
if (!acc_is_present (c, (N * sizeof (float))))
abort ();
acc_copyout (b, N * sizeof (float));
d = (float *) acc_deviceptr (b);
acc_memcpy_from_device (b, d, N * sizeof (float));
for (i = 0; i < N; i++)
{
@ -534,6 +536,12 @@ main (int argc, char **argv)
acc_free (d);
d = (float *) acc_deviceptr (b);
acc_unmap_data (b);
acc_free (d);
d = (float *) acc_deviceptr (c);
acc_unmap_data (c);

179
libgomp/testsuite/libgomp.oacc-c-c++-common/pr92843-1.c Normal file
View File

@ -0,0 +1,179 @@
/* 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;
}