Hi,
This patch provides a fix for PR92888, wherein global variables mapped
using an OpenACC 'declare' directive would not be visible to
device-pointer lookups.
Tested with offloading to nvptx. OK?
Thanks,
Julian
ChangeLog
2019-12-12 Julian Brown <jul...@codesourcery.com>
PR libgomp/92888
libgomp/
* oacc-parallel.c (GOACC_parallel_keyed): Add tgt_start in target
function address calculation.
* target.c (gomp_load_image_to_device): Record address range for
target_mem_desc for mapped functions and global variables, and adjust
tgt_offsets to be within that range.
(gomp_get_target_fn_addr): Add tgt_start in target function address
calculation.
* testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c:
Remove XFAIL.
commit 16e774d2ce86af90ff282b9126cf615e66e7efae
Author: Julian Brown <jul...@codesourcery.com>
Date: Mon Dec 9 11:04:58 2019 -0800
Find address range for offloaded functions and global variables (PR92888)
PR libgomp/92888
libgomp/
* oacc-parallel.c (GOACC_parallel_keyed): Add tgt_start in target
function address calculation.
* target.c (gomp_load_image_to_device): Record address range for
target_mem_desc for mapped functions and global variables, and adjust
offsets to be within that range.
(gomp_get_target_fn_addr): Add tgt_start in target function address
calculation.
* testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c:
Remove XFAIL.
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index f5ef5050bbd..5a5697cf6e6 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -377,7 +377,7 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *),
if (tgt_fn_key == NULL)
gomp_fatal ("target function wasn't mapped");
- tgt_fn = (void (*)) tgt_fn_key->tgt_offset;
+ tgt_fn = (void (*)) (tgt_fn_key->tgt->tgt_start + tgt_fn_key->tgt_offset);
}
else
tgt_fn = (void (*)) fn;
diff --git a/libgomp/target.c b/libgomp/target.c
index bb392dd1c8f..b023e3daf1a 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1759,6 +1759,8 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
tgt->device_descr = devicep;
splay_tree_node array = tgt->array;
+ uintptr_t max_addr = 0, min_addr = ~(uintptr_t) 0;
+
for (i = 0; i < num_funcs; i++)
{
splay_tree_key k = &array->key;
@@ -1766,6 +1768,10 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->host_end = k->host_start + 1;
k->tgt = tgt;
k->tgt_offset = target_table[i].start;
+ if (target_table[i].start < min_addr)
+ min_addr = target_table[i].start;
+ if (target_table[i].end > max_addr)
+ max_addr = target_table[i].end;
k->refcount = REFCOUNT_INFINITY;
k->virtual_refcount = 0;
k->aux = NULL;
@@ -1799,6 +1805,10 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
= k->host_start + (size_mask & (uintptr_t) host_var_table[i * 2 + 1]);
k->tgt = tgt;
k->tgt_offset = target_var->start;
+ if (target_var->start < min_addr)
+ min_addr = target_var->start;
+ if (target_var->end > max_addr)
+ max_addr = target_var->end;
k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY;
k->virtual_refcount = 0;
k->aux = NULL;
@@ -1808,6 +1818,17 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
array++;
}
+ /* Make the tgt_mem_desc cover all of the functions and variables so that
+ oacc-mem.c:lookup_dev can find mapped global variables properly. */
+ tgt->tgt_start = min_addr;
+ tgt->tgt_end = max_addr;
+
+ for (array = tgt->array, i = 0; i < num_vars + num_funcs; i++, array++)
+ {
+ splay_tree_key k = &array->key;
+ k->tgt_offset -= min_addr;
+ }
+
free (target_table);
}
@@ -2170,7 +2191,7 @@ gomp_get_target_fn_addr (struct gomp_device_descr *devicep,
if (tgt_fn == NULL)
return NULL;
- return (void *) tgt_fn->tgt_offset;
+ return (void *) (tgt_fn->tgt->tgt_start + tgt_fn->tgt_offset);
}
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c
index 7cd2936219a..0807bc9d694 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/acc_map_data-device_already-3.c
@@ -24,5 +24,5 @@ main ()
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "device address \\\[\[0-9a-fA-FxX\]+, \\\+8\\\] is already mapped" { xfail *-*-* } } TODO */
-/* { dg-shouldfail "TODO" { INV-AL-ID } } */
+/* { dg-output "device address \\\[\[0-9a-fA-FxX\]+, \\\+8\\\] is already mapped" } */
+/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92888-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92888-1.c
new file mode 100644
index 00000000000..0cd7f13656c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr92888-1.c
@@ -0,0 +1,19 @@
+/* Make sure that we can resolve back via 'acc_hostptr' an 'acc_deviceptr'
+ retrieved for a '#pragma acc declare'd variable. */
+
+#include <assert.h>
+#include <openacc.h>
+
+double global_var;
+#pragma acc declare create (global_var)
+
+int
+main ()
+{
+ void *global_var_p_d = acc_deviceptr (&global_var);
+ assert (acc_hostptr (global_var_p_d) == &global_var);
+
+ return 0;
+}
+
+/* { dg-xfail-run-if "PR92888" { ! openacc_host_selected } } */
