OpenACC device-pointer lookup with globally-mapped variables (PR92888)

commit 16e774d2ce86af90ff282b9126cf615e66e7efae
Author: Julian Brown <julian@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 } } */