libgomp: Fix declare target link with offset array-section mapping [PR116107]

Assume that 'int var[100]' is 'omp declare target link(var)'. When now
mapping an array section with offset such as 'map(to:var[20:10])',
the device-side link pointer has to store &<device-storage-data>[0] minus
the offset such that var[20] will access <device-storage-data>[0]. But
the offset calculation was missed such that the device-side 'var' pointed
to the first element of the mapped data - and var[20] points beyond at
some invalid memory.

	PR middle-end/116107

libgomp/ChangeLog:

	* target.c (gomp_map_vars_internal): Honor array mapping offsets
	with declare-target 'link' variables.
	* testsuite/libgomp.c-c++-common/target-link-2.c: New test.

2 files changed, 64 insertions, 2 deletions

diff --git a/libgomp/target.c b/libgomp/target.c
index aa01c13..efed6ad 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1820,8 +1820,11 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 		if (k->aux && k->aux->link_key)
 		  {
 		    /* Set link pointer on target to the device address of the
-		       mapped object.  */
-		    void *tgt_addr = (void *) (tgt->tgt_start + k->tgt_offset);
+		       mapped object.  Also deal with offsets due to
+		       array-section mapping.  */
+		    void *tgt_addr = (void *) (tgt->tgt_start + k->tgt_offset
+					       - (k->host_start
+						  - k->aux->link_key->host_start));
 		    /* We intentionally do not use coalescing here, as it's not
 		       data allocated by the current call to this function.  */
 		    gomp_copy_host2dev (devicep, aq, (void *) n->tgt_offset,
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-link-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-link-2.c
new file mode 100644
index 0000000..15da165
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-link-2.c
@@ -0,0 +1,59 @@
+/* { dg-do run }  */
+/* PR middle-end/116107  */
+
+#include <omp.h>
+
+int arr[15] = {1,2,3,4,5,6,7,8,9,10,11,12,13,14,15};
+#pragma omp declare target link(arr)
+
+#pragma omp begin declare target
+void f(int *res)
+{
+  __builtin_memcpy (res, &arr[5], sizeof(int)*10);
+}
+
+void g(int *res)
+{
+  __builtin_memcpy (res, &arr[3], sizeof(int)*10);
+}
+#pragma omp end declare target
+
+int main()
+{
+  int res[10], res2;
+  for (int dev = 0; dev < omp_get_num_devices(); dev++)
+    {
+      __builtin_memset (res, 0, sizeof (res));
+      res2 = 99;
+
+      #pragma omp target enter data map(arr[5:10]) device(dev)
+
+      #pragma omp target map(from: res) device(dev)
+	f (res);
+
+      #pragma omp target map(from: res2) device(dev)
+	res2 = arr[5];
+
+      if (res2 != 6)
+	__builtin_abort ();
+      for (int i = 0; i < 10; i++)
+	if (res[i] != 6 + i)
+	  __builtin_abort ();
+
+      #pragma omp target exit data map(release:arr[5:10]) device(dev)
+
+      for (int i = 0; i < 15; i++)
+	arr[i] *= 10;
+      __builtin_memset (res, 0, sizeof (res));
+
+      #pragma omp target enter data map(arr[3:10]) device(dev)
+
+      #pragma omp target map(from: res) device(dev)
+	g (res);
+
+      for (int i = 0; i < 10; i++)
+	if (res[i] != (4 + i)*10)
+	  __builtin_abort ();
+    }
+  return 0;
+}