mirror/gcc
1
0
Fork 0
mirror of https://github.com/gcc-mirror/gcc.git synced 2024-11-21 13:40:47 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

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

Browse Source 
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.
This commit is contained in:
Tobias Burnus 2024-07-29 11:40:38 +02:00
parent b3176b620f
commit 14c47e7eb0
2 changed files with 64 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
libgomp/target.c
View File

@ -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,

59
libgomp/testsuite/libgomp.c-c++-common/target-link-2.c Normal file
View File

@ -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;
}