mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-02-19 20:30:28 +08:00
Code Issues Packages Projects Releases Wiki Activity

[PR92116, PR92877] [OpenACC] Replace 'openacc.data_environ' by standard libgomp mechanics

Browse Source 
libgomp/
	PR libgomp/92116
	PR libgomp/92877
	* oacc-mem.c (lookup_dev): Reimplement.  Adjust all users.
	* libgomp.h (struct acc_dispatch_t): Remove 'data_environ' member.
	Adjust all users.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c:
	Remove XFAIL.
	* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c:
	Likewise.
	* testsuite/libgomp.oacc-c-c++-common/pr92877-1.c: New file.

Co-Authored-By: Julian Brown <julian@codesourcery.com>

From-SVN: r279147
This commit is contained in:
Thomas Schwinge 2019-12-09 23:52:56 +01:00 committed by Thomas Schwinge
parent cec41816c1
commit 47afc7b4dd
8 changed files with 69 additions and 117 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
libgomp/ChangeLog
View File

@ -1,3 +1,18 @@
2019-12-09 Thomas Schwinge <thomas@codesourcery.com>
Julian Brown <julian@codesourcery.com>
PR libgomp/92116
PR libgomp/92877
* oacc-mem.c (lookup_dev): Reimplement. Adjust all users.
* libgomp.h (struct acc_dispatch_t): Remove 'data_environ' member.
Adjust all users.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c:
Remove XFAIL.
* testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr92877-1.c: New file.
2019-12-09 Thomas Schwinge <thomas@codesourcery.com>
PR libgomp/92503

10
libgomp/libgomp.h
View File

@ -1025,13 +1025,6 @@ splay_compare (splay_tree_key x, splay_tree_key y)
typedef struct acc_dispatch_t
{
/* This is a linked list of data mapped using the
acc_map_data/acc_unmap_data or "acc enter data"/"acc exit data" pragmas.
Unlike mapped_data in the goacc_thread struct, unmapping can
happen out-of-order with respect to mapping. */
/* This is guarded by the lock in the "outer" struct gomp_device_descr. */
struct target_mem_desc *data_environ;
/* Execute. */
__typeof (GOMP_OFFLOAD_openacc_exec) *exec_func;
@ -1132,8 +1125,7 @@ struct gomp_device_descr
enum gomp_device_state state;
/* OpenACC-specific data and functions. */
/* This is mutable because of its mutable data_environ and target_data
members. */
/* This is mutable because of its mutable target_data member. */
acc_dispatch_t openacc;
};

2
libgomp/oacc-host.c
View File

@ -264,8 +264,6 @@ static struct gomp_device_descr host_dispatch =
.state = GOMP_DEVICE_UNINITIALIZED,
.openacc = {
.data_environ = NULL,
.exec_func = host_openacc_exec,
.create_thread_data_func = host_openacc_create_thread_data,

131
libgomp/oacc-mem.c
View File

@ -50,44 +50,42 @@ lookup_host (struct gomp_device_descr *dev, void *h, size_t s)
return key;
}
/* Return block containing [D->S), or NULL if not contained.
The list isn't ordered by device address, so we have to iterate
over the whole array. This is not expected to be a common
operation. The device lock associated with TGT must be locked on entry, and
remains locked on exit. */
/* Helper for lookup_dev. Iterate over splay tree. */
static splay_tree_key
lookup_dev (struct target_mem_desc *tgt, void *d, size_t s)
lookup_dev_1 (splay_tree_node node, uintptr_t d, size_t s)
{
int i;
struct target_mem_desc *t;
splay_tree_key key = &node->key;
if (d >= key->tgt->tgt_start && d + s <= key->tgt->tgt_end)
return key;
if (!tgt)
return NULL;
key = NULL;
if (node->left)
key = lookup_dev_1 (node->left, d, s);
if (!key && node->right)
key = lookup_dev_1 (node->right, d, s);
for (t = tgt; t != NULL; t = t->prev)
{
if (t->tgt_start <= (uintptr_t) d && t->tgt_end >= (uintptr_t) d + s)
break;
}
if (!t)
return NULL;
for (i = 0; i < t->list_count; i++)
{
void * offset;
splay_tree_key k = &t->array[i].key;
offset = d - t->tgt_start + k->tgt_offset;
if (k->host_start + offset <= (void *) k->host_end)
return k;
}
return NULL;
return key;
}
/* Return block containing [D->S), or NULL if not contained.
This iterates over the splay tree. This is not expected to be a common
operation.
The device lock associated with MEM_MAP must be locked on entry, and remains
locked on exit. */
static splay_tree_key
lookup_dev (splay_tree mem_map, void *d, size_t s)
{
if (!mem_map || !mem_map->root)
return NULL;
return lookup_dev_1 (mem_map->root, (uintptr_t) d, s);
}
/* OpenACC is silent on how memory exhaustion is indicated. We return
NULL. */
@ -147,7 +145,7 @@ acc_free (void *d)
/* We don't have to call lazy open here, as the ptr value must have
been returned by acc_malloc. It's not permitted to pass NULL in
(unless you got that null from acc_malloc). */
if ((k = lookup_dev (acc_dev->openacc.data_environ, d, 1)))
if ((k = lookup_dev (&acc_dev->mem_map, d, 1)))
{
void *offset = d - k->tgt->tgt_start + k->tgt_offset;
void *h = k->host_start + offset;
@ -300,7 +298,7 @@ acc_hostptr (void *d)
gomp_mutex_lock (&acc_dev->lock);
n = lookup_dev (acc_dev->openacc.data_environ, d, 1);
n = lookup_dev (&acc_dev->mem_map, d, 1);
if (!n)
{
@ -395,7 +393,7 @@ acc_map_data (void *h, void *d, size_t s)
(int)s);
}
if (lookup_dev (thr->dev->openacc.data_environ, d, s))
if (lookup_dev (&thr->dev->mem_map, d, s))
{
gomp_mutex_unlock (&acc_dev->lock);
gomp_fatal ("device address [%p, +%d] is already mapped", (void *)d,
@ -418,11 +416,6 @@ acc_map_data (void *h, void *d, size_t s)
thr->api_info = NULL;
}
}
gomp_mutex_lock (&acc_dev->lock);
tgt->prev = acc_dev->openacc.data_environ;
acc_dev->openacc.data_environ = tgt;
gomp_mutex_unlock (&acc_dev->lock);
}
void
@ -482,25 +475,11 @@ acc_unmap_data (void *h)
if (t->refcount == 2)
{
struct target_mem_desc *tp;
/* This is the last reference, so pull the descriptor off the
chain. This avoids gomp_unmap_vars via gomp_unmap_tgt from
freeing the device memory. */
t->tgt_end = 0;
t->to_free = 0;
for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
tp = t, t = t->prev)
if (n->tgt == t)
{
if (tp)
tp->prev = t->prev;
else
acc_dev->openacc.data_environ = t->prev;
break;
}
}
gomp_mutex_unlock (&acc_dev->lock);
@ -597,13 +576,7 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
/* Initialize dynamic refcount. */
tgt->list[0].key->dynamic_refcount = 1;
gomp_mutex_lock (&acc_dev->lock);
d = tgt->to_free;
tgt->prev = acc_dev->openacc.data_environ;
acc_dev->openacc.data_environ = tgt;
gomp_mutex_unlock (&acc_dev->lock);
}
if (profiling_p)
@ -749,21 +722,6 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
if (n->refcount == 0)
{
if (n->tgt->refcount == 2)
{
struct target_mem_desc *tp, *t;
for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
tp = t, t = t->prev)
if (n->tgt == t)
{
if (tp)
tp->prev = t->prev;
else
acc_dev->openacc.data_environ = t->prev;
break;
}
}
if (f & FLAG_COPYOUT)
{
goacc_aq aq = get_goacc_asyncqueue (async);
@ -954,11 +912,6 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
/* Initialize dynamic refcount. */
tgt->list[0].key->dynamic_refcount = 1;
gomp_mutex_lock (&acc_dev->lock);
tgt->prev = acc_dev->openacc.data_environ;
acc_dev->openacc.data_environ = tgt;
gomp_mutex_unlock (&acc_dev->lock);
}
void
@ -1009,26 +962,6 @@ gomp_acc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
if (n->refcount == 0)
{
if (t->refcount == minrefs)
{
/* This is the last reference, so pull the descriptor off the
chain. This prevents gomp_unmap_vars via gomp_unmap_tgt from
freeing the device memory. */
struct target_mem_desc *tp;
for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
tp = t, t = t->prev)
{
if (n->tgt == t)
{
if (tp)
tp->prev = t->prev;
else
acc_dev->openacc.data_environ = t->prev;
break;
}
}
}
/* Set refcount to 1 to allow gomp_unmap_vars to unmap it. */
n->refcount = 1;
t->refcount = minrefs;

1
libgomp/target.c
View File

@ -2897,7 +2897,6 @@ gomp_target_init (void)
current_device.type = current_device.get_type_func ();
current_device.mem_map.root = NULL;
current_device.state = GOMP_DEVICE_UNINITIALIZED;
current_device.openacc.data_environ = NULL;
for (i = 0; i < new_num_devices; i++)
{
current_device.target_id = i;

4
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4-2.c
View File

@ -25,7 +25,5 @@ main ()
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
TODO PR92877
{ dg-output "libgomp: cuMemGetAddressRange_v2 error: named symbol not found" { target openacc_nvidia_accel_selected } }
{ dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+64\\\]" { xfail *-*-* } }
{ dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+64\\\]" }
{ dg-shouldfail "" } */

4
libgomp/testsuite/libgomp.oacc-c-c++-common/acc_free-pr92503-4.c
View File

@ -26,7 +26,5 @@ main ()
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
TODO PR92877
{ dg-output "libgomp: cuMemGetAddressRange_v2 error: named symbol not found" { target openacc_nvidia_accel_selected } }
{ dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+87\\\]" { xfail *-*-* } }
{ dg-output "refusing to free device memory space at \[0-9a-fA-FxX\]+ that is still mapped at \\\[\[0-9a-fA-FxX\]+,\\\+87\\\]" }
{ dg-shouldfail "" } */

19
libgomp/testsuite/libgomp.oacc-c-c++-common/pr92877-1.c Normal file
View File

@ -0,0 +1,19 @@
/* Make sure that we can resolve back via 'acc_hostptr' an 'acc_deviceptr'
retrieved for a structured mapping. */
#include <assert.h>
#include <openacc.h>
int
main ()
{
int var;
#pragma acc data create (var)
{
void *var_p_d = acc_deviceptr (&var);
assert (acc_hostptr (var_p_d) == &var);
}
return 0;
}