b25ea7ab78
For C/C++ pointers, default implicit mapping firstprivatizes the pointer but if the memory it points to is mapped, the it is updated to point to the device memory (by attaching a zero sized array section of the pointed-to storage). However, if the pointed-to storage wasn't mapped, the pointer was set to NULL on the device side (OpenMP 5.0/5.1 semantic). With this commit, the pointer retains the on-host address in that case (OpenMP 5.2 semantic). The new semantic avoids an explicit map/firstprivate/is_device_ptr in the following sensible cases: Special values (e.g. pointer or 0x1, 0x2 etc.), explicitly device allocated memory (e.g. omp_target_alloc), and with (unified) shared memory. (Note: With (U)SM, mappings still must be tracked, at least when omp_target_associate_ptr does not fail when passing in two destinct pointers.) libgomp/ PR middle-end/110270 * target.c (gomp_map_vars_internal): Copy host value instead of NULL for GOMP_MAP_ZERO_LEN_ARRAY_SECTION if not mapped. * libgomp.texi (OpenMP 5.2 Impl.): Mark as 'Y'. * testsuite/libgomp.c/target-19.c: Update expected value. * testsuite/libgomp.c++/target-18.C: Likewise. * testsuite/libgomp.c++/target-19.C: Likewise. * testsuite/libgomp.c-c++-common/requires-unified-addr-2.c: New test. * testsuite/libgomp.c-c++-common/target-implicit-map-3.c: New test. * testsuite/libgomp.c-c++-common/target-implicit-map-4.c: New test.
85 lines
2.4 KiB
C
85 lines
2.4 KiB
C
/* PR middle-end/110270 */
|
|
|
|
/* OpenMP 5.2's 'defaultmap(default : pointer) for C/C++ pointers retains the
|
|
pointer value instead of setting it to NULL if the pointer cannot be found.
|
|
Contrary to requires-unified-addr-1.c which is valid OpenMP 5.0/5.1/5.2,
|
|
this testcase is only valid since OpenMP 5.2. */
|
|
|
|
/* This is kind of a follow-up to the requires-unified-addr-1.c testcase
|
|
and PR libgomp/109837 */
|
|
|
|
|
|
#include <assert.h>
|
|
#include <omp.h>
|
|
#include <stdio.h>
|
|
#include <stdlib.h>
|
|
|
|
#pragma omp requires unified_address
|
|
|
|
#define N 15
|
|
|
|
void
|
|
test_device (int dev)
|
|
{
|
|
struct st {
|
|
int *ptr;
|
|
int n;
|
|
};
|
|
struct st s;
|
|
|
|
s.n = 10;
|
|
s.ptr = (int *) omp_target_alloc (sizeof (int)*s.n, dev);
|
|
int *ptr1 = (int *) omp_target_alloc (sizeof (int)*N, dev);
|
|
assert (s.ptr != NULL);
|
|
assert (ptr1 != NULL);
|
|
|
|
int q[4] = {1,2,3,4};
|
|
int *qptr;
|
|
#pragma omp target enter data map(q) device(device_num: dev)
|
|
#pragma omp target data use_device_addr(q) device(device_num: dev)
|
|
qptr = q;
|
|
|
|
#pragma omp target map(to:s) device(device_num: dev)
|
|
for (int i = 0; i < s.n; i++)
|
|
s.ptr[i] = 23*i;
|
|
|
|
int *ptr2 = &s.ptr[3];
|
|
|
|
/* s.ptr is not mapped (but omp_target_alloc'ed) thus ptr2 shall retain its value. */
|
|
#pragma omp target device(device_num: dev) /* implied: defaultmap(default : pointer) */
|
|
for (int i = 0; i < 4; i++)
|
|
*(qptr++) = ptr2[i];
|
|
|
|
#pragma omp target exit data map(q) device(device_num: dev)
|
|
for (int i = 0; i < 4; i++)
|
|
q[i] = 23 * (i+3);
|
|
|
|
/* ptr1 retains the value as it is not mapped (but it is omp_target_alloc'ed). */
|
|
#pragma omp target defaultmap(default : pointer) device(device_num: dev)
|
|
for (int i = 0; i < N; i++)
|
|
ptr1[i] = 11*i;
|
|
|
|
int *ptr3 = (int *) malloc (sizeof (int)*N);
|
|
assert (0 == omp_target_memcpy(ptr3, ptr1, N * sizeof(int), 0, 0,
|
|
omp_get_initial_device(), dev));
|
|
for (int i = 0; i < N; i++)
|
|
assert (ptr3[i] == 11*i);
|
|
|
|
free (ptr3);
|
|
omp_target_free (ptr1, dev);
|
|
omp_target_free (s.ptr, dev);
|
|
}
|
|
|
|
int
|
|
main()
|
|
{
|
|
int ntgts = omp_get_num_devices();
|
|
if (ntgts)
|
|
fprintf (stderr, "Offloading devices exist\n"); /* { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } */
|
|
else
|
|
fprintf (stderr, "Only host fallback\n"); /* { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } */
|
|
|
|
for (int i = 0; i <= ntgts; i++)
|
|
test_device (i);
|
|
return 0;
|
|
}
|