https://gcc.gnu.org/g:c9e52a1a3d2c2970065c254a414bab76f798ce7d
commit c9e52a1a3d2c2970065c254a414bab76f798ce7d
Author: Tobias Burnus <tbur...@baylibre.com>
Date: Mon Jul 29 12:50:56 2024 +0200
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.
(cherry picked from commit 14c47e7eb06e8b95913794f6059560fc2fa6de91)
Diff:
---
libgomp/ChangeLog | 11 ++++
libgomp/target.c | 7 ++-
.../testsuite/libgomp.c-c++-common/target-link-2.c | 59 ++++++++++++++++++++++
3 files changed, 75 insertions(+), 2 deletions(-)
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 555f1f126f26..aaa2e8defe3e 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,14 @@
+------ libgomp/ChangeLog ------
+2024-07-29 Tobias Burnus <tbur...@baylibre.com>
+
+ Backported from master:
+ 2024-07-29 Tobias Burnus <tbur...@baylibre.com>
+
+ PR middle-end/116107
+ * 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.
+
2024-05-07 Jakub Jelinek <ja...@redhat.com>
Backported from master:
diff --git a/libgomp/target.c b/libgomp/target.c
index caa501c27acb..eb02d478e109 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1892,8 +1892,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 000000000000..15da1656ebf9
--- /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;
+}
