mirrors/gcc
1
0
Fork 0
mirror of https://gcc.gnu.org/git/gcc.git synced 2024-11-23 02:44:18 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

libgomp: with USM, init 'link' variables with host address

Browse Source 
If requires unified_shared_memory or self_maps is set, make
'declare target link' variables to point initially to the host pointer.

libgomp/ChangeLog:

	* target.c (gomp_load_image_to_device): For requires
	unified_shared_memory, update 'link' vars to point to the host var.
	* testsuite/libgomp.c-c++-common/target-link-3.c: New test.
	* testsuite/libgomp.c-c++-common/target-link-4.c: New test.
This commit is contained in:
Tobias Burnus 2024-09-24 17:41:39 +02:00
parent 79a3d3da8c
commit 4cb20dc043
3 changed files with 110 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
libgomp/target.c
View File

@ -2454,6 +2454,12 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
array++;
if (is_link_var
&& (omp_requires_mask
& (GOMP_REQUIRES_UNIFIED_SHARED_MEMORY | GOMP_REQUIRES_SELF_MAPS)))
gomp_copy_host2dev (devicep, NULL, (void *) target_var->start,
&k->host_start, sizeof (void *), false, NULL);
}
/* Last entry is for the ICV struct variable; if absent, start = end = 0. */

52
libgomp/testsuite/libgomp.c-c++-common/target-link-3.c Normal file
View File

@ -0,0 +1,52 @@
/* { dg-do run } */
#include <stdint.h>
#include <omp.h>
#pragma omp requires unified_shared_memory
int A[3] = {-3,-4,-5};
static int q = -401;
#pragma omp declare target link(A, q)
#pragma omp begin declare target
void
f (uintptr_t *pA, uintptr_t *pq)
{
if (A[0] != 1 || A[1] != 2 || A[2] != 3 || q != 42)
__builtin_abort ();
A[0] = 13;
A[1] = 14;
A[2] = 15;
q = 23;
*pA = (uintptr_t) &A[0];
*pq = (uintptr_t) &q;
}
#pragma omp end declare target
int
main ()
{
uintptr_t hpA = (uintptr_t) &A[0];
uintptr_t hpq = (uintptr_t) &q;
uintptr_t dpA, dpq;
A[0] = 1;
A[1] = 2;
A[2] = 3;
q = 42;
for (int i = 0; i <= omp_get_num_devices (); ++i)
{
#pragma omp target device(device_num: i) map(dpA, dpq)
f (&dpA, &dpq);
if (hpA != dpA || hpq != dpq)
__builtin_abort ();
if (A[0] != 13 || A[1] != 14 || A[2] != 15 || q != 23)
__builtin_abort ();
A[0] = 1;
A[1] = 2;
A[2] = 3;
q = 42;
}
}

52
libgomp/testsuite/libgomp.c-c++-common/target-link-4.c Normal file
View File

@ -0,0 +1,52 @@
/* { dg-do run } */
#include <stdint.h>
#include <omp.h>
#pragma omp requires self_maps
int A[3] = {-3,-4,-5};
static int q = -401;
#pragma omp declare target link(A, q)
#pragma omp begin declare target
void
f (uintptr_t *pA, uintptr_t *pq)
{
if (A[0] != 1 || A[1] != 2 || A[2] != 3 || q != 42)
__builtin_abort ();
A[0] = 13;
A[1] = 14;
A[2] = 15;
q = 23;
*pA = (uintptr_t) &A[0];
*pq = (uintptr_t) &q;
}
#pragma omp end declare target
int
main ()
{
uintptr_t hpA = (uintptr_t) &A[0];
uintptr_t hpq = (uintptr_t) &q;
uintptr_t dpA, dpq;
A[0] = 1;
A[1] = 2;
A[2] = 3;
q = 42;
for (int i = 0; i <= omp_get_num_devices (); ++i)
{
#pragma omp target device(device_num: i) map(dpA, dpq)
f (&dpA, &dpq);
if (hpA != dpA || hpq != dpq)
__builtin_abort ();
if (A[0] != 13 || A[1] != 14 || A[2] != 15 || q != 23)
__builtin_abort ();
A[0] = 1;
A[1] = 2;
A[2] = 3;
q = 42;
}
}