The main idea of 'link' is to permit putting only a subset of a
huge array on the device. Well, in order to make this work properly,
it requires that one can map an array section, which does not
start with the first element.
This patch adjusts the pointers such, that this actually works.
(Tested on x86-64-gnu-linux with Nvptx offloading.)
Comments, suggestions, remarks before I commit it?
Tobias
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.
libgomp/target.c | 7 ++-
.../testsuite/libgomp.c-c++-common/target-link-2.c | 59 ++++++++++++++++++++++
2 files changed, 64 insertions(+), 2 deletions(-)
diff --git a/libgomp/target.c b/libgomp/target.c
index aa01c1367b9..e3e648f5443 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 00000000000..4ff4080da76
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-link-2.c
@@ -0,0 +1,59 @@
+/* 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++)
+ res[i] *= 10;
+ __builtin_abort ();
+
+ #pragma omp target enter data map(arr[3:10]) device(dev)
+ __builtin_memset (res, 0, sizeof (res));
+
+ #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;
+}
