shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

[OpenMP] Preserve the original address when use_device_ptr/addr lookup fails. (#174659)

Browse Source 
As per OpenMP 5.1, we need to assume that when the lookup for
`use_device_ptr/addr` fails, the incoming pointer was already device
accessible.
    
Prior to 5.1, a lookup-failure meant a user-error (for
`use_device_ptr`),
so we could do anything in that scenario. For `use_device_addr`,
it was always incorrect to set the address to null.

OpenMP 6.1 adds a way to retain the previous behavior of nullifying a
pointer
when the lookup fails. That will be tackled by the PR stack
starting with https://github.com/llvm/llvm-project/pull/169603.
This commit is contained in:
Abhinav Gaba 2026-01-09 14:42:30 -08:00 committed by GitHub
parent acb78bde6f
commit cd81aae57a
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
13 changed files with 78 additions and 118 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
clang/docs/OpenMPSupport.rst
View File

@ -266,6 +266,8 @@ implementation.
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | has_device_addr clause on target construct | :none:`unclaimed` | |
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | use_device_ptr/addr preserve host address when lookup fails | :good:`done` | https://github.com/llvm/llvm-project/pull/174659 |
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | iterators in map clause or motion clauses | :none:`done` | https://github.com/llvm/llvm-project/pull/159112 |
+------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
| device | indirect clause on declare target directive | :part:`In Progress` | |

2
clang/docs/ReleaseNotes.rst
View File

@ -988,6 +988,8 @@ OpenMP Support
- Added parsing and semantic analysis support for ``need_device_ptr`` modifier
to accept an optional fallback argument (``fb_nullify`` or ``fb_preserve``)
with OpenMP >= 61.
- ``use_device_ptr`` and ``use_device_addr`` now preserve the original host
address when lookup fails.
Improvements
^^^^^^^^^^^^

36
offload/libomptarget/omptarget.cpp
View File

@ -683,9 +683,39 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
<< " new";
if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
ODBG(ODT_Mapping) << "Returning device pointer " << TgtPtrBase;
uintptr_t Delta = reinterpret_cast<uintptr_t>(HstPtrBegin) -
reinterpret_cast<uintptr_t>(HstPtrBase);
void *TgtPtrBase;
if (TgtPtrBegin) {
// Lookup succeeded, return device pointer adjusted by delta
TgtPtrBase = reinterpret_cast<void *>(
reinterpret_cast<uintptr_t>(TgtPtrBegin) - Delta);
ODBG(ODT_Mapping) << "Returning device pointer " << TgtPtrBase;
} else {
// Lookup failed. So we have to decide what to do based on the
// requested fallback behavior.
//
// Treat "preserve" as the default fallback behavior, since as per
// OpenMP 5.1, for use_device_ptr/addr, when there's no corresponding
// device pointer to translate into, it's the user's responsibility to
// ensure that the host address is device-accessible.
//
// OpenMP 5.1, sec 2.14.2, target data construct, p 188, l26-31:
// If a list item that appears in a use_device_ptr clause ... does not
// point to a mapped object, it must contain a valid device address for
// the target device, and the list item references are instead converted
// to references to a local device pointer that refers to this device
// address.
//
// TODO: Add a new map-type bit to support OpenMP 6.1's `fb_nullify`
// and set the result to `nullptr - Delta`. Note that `fb_nullify` is
// already the default for `need_device_ptr`, but clang/flang do not
// support its codegen yet.
TgtPtrBase = reinterpret_cast<void *>(
reinterpret_cast<uintptr_t>(HstPtrBegin) - Delta);
ODBG(ODT_Mapping) << "Returning host pointer " << TgtPtrBase
<< " as fallback (lookup failed)";
}
ArgsBase[I] = TgtPtrBase;
}

2
offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
View File

@ -7,8 +7,6 @@
// list-item is device-accessible, even if it was not
// previously mapped.
// XFAIL: *
#include <stdio.h>
int h[10];
int *ph = &h[0];

20
offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
View File

@ -6,15 +6,6 @@
// Test for various cases of use_device_addr on an array-section.
// The corresponding data is not previously mapped.
// Note that this tests for the current behavior wherein if a lookup fails,
// the runtime returns nullptr, instead of the original host-address.
// That was compatible with OpenMP 5.0, where it was a user error if
// corresponding storage didn't exist, but with 5.1+, the runtime needs to
// return the host address, as it needs to assume that the host-address is
// device-accessible, as the user has guaranteed it.
// Once the runtime returns the original host-address when the lookup fails, the
// test will need to be updated.
int g, h[10];
int *ph = &h[0];
@ -34,7 +25,7 @@ struct S {
int *mapped_ptr_ph3 =
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (B) use_device_addr/map: different operands, same base-pointer.
@ -56,7 +47,7 @@ struct S {
int *mapped_ptr_ph3 =
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (D) use_device_addr/map: one of two maps with matching base-pointer.
@ -78,8 +69,7 @@ struct S {
int **mapped_ptr_paa02 =
(int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
mapped_ptr_paa02 != original_paa02,
&paa[0][2] == (int **)nullptr + 2);
mapped_ptr_paa02 != original_paa02, &paa[0][2] == original_paa02);
}
// (F) use_device_addr/map: different operands, same base-array.
@ -108,7 +98,7 @@ struct S {
}
int *original_paa020 = &paa[0][2][0];
int **original_paa0 = (int **)&paa[0];
void *original_paa0 = &paa[0];
// (H) use_device_addr/map: different base-pointers.
// No corresponding storage for use_device_addr opnd, lookup should fail.
@ -120,7 +110,7 @@ struct S {
int **mapped_ptr_paa0 =
(int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
mapped_ptr_paa0 == nullptr, &paa[0] == original_paa0);
}
// (I) use_device_addr/map: one map with different, one with same base-ptr.

28
offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
View File

@ -6,15 +6,6 @@
// Test for various cases of use_device_addr on an array-section on a reference.
// The corresponding data is not previously mapped.
// Note that this tests for the current behavior wherein if a lookup fails,
// the runtime returns nullptr, instead of the original host-address.
// That was compatible with OpenMP 5.0, where it was a user error if
// corresponding storage didn't exist, but with 5.1+, the runtime needs to
// return the host address, as it needs to assume that the host-address is
// device-accessible, as the user has guaranteed it.
// Once the runtime returns the original host-address when the lookup fails, the
// test will need to be updated.
int g_ptee;
int &g = g_ptee;
@ -35,15 +26,13 @@ struct S {
int **original_paa02 = &paa[0][2];
// (A) No corresponding map, lookup should fail.
// EXPECTED: A: 1 1 1
// CHECK: A: 1 1 0
// FIXME: ph is not being privatized in the region.
// CHECK: A: 1 1 1
#pragma omp target data use_device_addr(ph[3 : 4])
{
int *mapped_ptr_ph3 =
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (B) use_device_addr/map: different operands, same base-pointer.
@ -61,15 +50,13 @@ struct S {
// (C) use_device_addr/map: different base-pointers.
// No corresponding storage, lookup should fail.
// EXPECTED: C: 1 1 1
// CHECK: C: 1 1 0
// FIXME: ph is not being privatized in the region.
// CHECK: C: 1 1 1
#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
{
int *mapped_ptr_ph3 =
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (D) use_device_addr/map: one of two maps with matching base-pointer.
@ -93,8 +80,7 @@ struct S {
int **mapped_ptr_paa02 =
(int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
mapped_ptr_paa02 != original_paa02,
&paa[0][2] == (int **)nullptr + 2);
mapped_ptr_paa02 != original_paa02, &paa[0][2] == original_paa02);
}
// (F) use_device_addr/map: different operands, same base-array.
@ -123,7 +109,7 @@ struct S {
}
int *original_paa020 = &paa[0][2][0];
int **original_paa0 = (int **)&paa[0];
void *original_paa0 = &paa[0];
// (H) use_device_addr/map: different base-pointers.
// No corresponding storage for use_device_addr opnd, lookup should fail.
@ -135,7 +121,7 @@ struct S {
int **mapped_ptr_paa0 =
(int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
mapped_ptr_paa0 == nullptr, &paa[0] == original_paa0);
}
// (I) use_device_addr/map: one map with different, one with same base-ptr.

2
offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
View File

@ -7,8 +7,6 @@
// list-item is device-accessible, even if it was not
// previously mapped.
// XFAIL: *
#include <stdio.h>
int x;

21
offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
View File

@ -6,15 +6,6 @@
// Test for various cases of use_device_addr on a variable (not a section).
// The corresponding data is not previously mapped.
// Note that this tests for the current behavior wherein if a lookup fails,
// the runtime returns nullptr, instead of the original host-address.
// That was compatible with OpenMP 5.0, where it was a user error if
// corresponding storage didn't exist, but with 5.1+, the runtime needs to
// return the host address, as it needs to assume that the host-address is
// device-accessible, as the user has guaranteed it.
// Once the runtime returns the original host-address when the lookup fails, the
// test will need to be updated.
int g, h[10];
int *ph = &h[0];
@ -36,7 +27,7 @@ struct S {
void *mapped_ptr_g =
omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_g == nullptr,
mapped_ptr_g != original_addr_g, (void *)&g == nullptr);
mapped_ptr_g != original_addr_g, &g == original_addr_g);
}
// (B) Lookup should succeed.
@ -56,7 +47,7 @@ struct S {
void *mapped_ptr_h =
omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
printf("C: %d %d %d\n", mapped_ptr_h == nullptr,
mapped_ptr_h != original_addr_h, (void *)&h == nullptr);
mapped_ptr_h != original_addr_h, &h == original_addr_h);
}
// (D) Lookup should succeed.
@ -76,7 +67,7 @@ struct S {
void *mapped_ptr_ph =
omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_ph == nullptr,
mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (F) Lookup should succeed.
@ -97,7 +88,7 @@ struct S {
void *mapped_ptr_ph =
omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
printf("G: %d %d %d\n", mapped_ptr_ph == nullptr,
mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (H) Maps both pointee and pointer. Lookup for pointer should succeed.
@ -117,7 +108,7 @@ struct S {
void *mapped_ptr_paa =
omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
printf("I: %d %d %d\n", mapped_ptr_paa == nullptr,
mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (J) Maps pointee only, but use_device_addr operand is pointer.
@ -128,7 +119,7 @@ struct S {
void *mapped_ptr_paa =
omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
printf("J: %d %d %d\n", mapped_ptr_paa == nullptr,
mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (K) Lookup should succeed.

21
offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
View File

@ -6,15 +6,6 @@
// Test for various cases of use_device_addr on a reference variable.
// The corresponding data is not previously mapped.
// Note that this tests for the current behavior wherein if a lookup fails,
// the runtime returns nullptr, instead of the original host-address.
// That was compatible with OpenMP 5.0, where it was a user error if
// corresponding storage didn't exist, but with 5.1+, the runtime needs to
// return the host address, as it needs to assume that the host-address is
// device-accessible, as the user has guaranteed it.
// Once the runtime returns the original host-address when the lookup fails, the
// test will need to be updated.
int g_ptee;
int &g = g_ptee;
@ -43,7 +34,7 @@ struct S {
void *mapped_ptr_g =
omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_g == nullptr,
mapped_ptr_g != original_addr_g, (void *)&g == nullptr);
mapped_ptr_g != original_addr_g, &g == original_addr_g);
}
// (B) Lookup should succeed.
@ -63,7 +54,7 @@ struct S {
void *mapped_ptr_h =
omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
printf("C: %d %d %d\n", mapped_ptr_h == nullptr,
mapped_ptr_h != original_addr_h, (void *)&h == nullptr);
mapped_ptr_h != original_addr_h, &h == original_addr_h);
}
// (D) Lookup should succeed.
@ -83,7 +74,7 @@ struct S {
void *mapped_ptr_ph =
omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_ph == nullptr,
mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (F) Lookup should succeed.
@ -104,7 +95,7 @@ struct S {
void *mapped_ptr_ph =
omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
printf("G: %d %d %d\n", mapped_ptr_ph == nullptr,
mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (H) Maps both pointee and pointer. Lookup for pointer should succeed.
@ -124,7 +115,7 @@ struct S {
void *mapped_ptr_paa =
omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
printf("I: %d %d %d\n", mapped_ptr_paa == nullptr,
mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (J) Maps pointee only, but use_device_addr operand is pointer.
@ -135,7 +126,7 @@ struct S {
void *mapped_ptr_paa =
omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
printf("J: %d %d %d\n", mapped_ptr_paa == nullptr,
mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (K) Lookup should succeed.

5
offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
View File

@ -1,5 +1,5 @@
// RUN: %libomptarget-compile-generic -fopenmp-version=51 -g
// RUN: env LIBOMPTARGET_INFO=64 %libomptarget-run-fail-generic 2>&1 \
// RUN: env LIBOMPTARGET_INFO=64 %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
// FIXME: Fails due to optimized debugging in 'ptxas'
@ -20,7 +20,8 @@ int main() {
// counterpart
#pragma omp target data use_device_addr(x)
{
// CHECK-NOT: device addr=0x[[#%x,HOST_ADDR:]]
// Even when the lookup fails, x should retain its host address.
// CHECK: device addr=0x[[#HOST_ADDR]]
fprintf(stderr, "device addr=%p\n", x);
}
}

19
offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
View File

@ -6,15 +6,6 @@
// Test for various cases of use_device_ptr on a variable.
// The corresponding data is not previously mapped.
// Note that this tests for the current behavior wherein if a lookup fails,
// the runtime returns nullptr, instead of the original host-address.
// That was compatible with OpenMP 5.0, where it was a user error if
// corresponding storage didn't exist, but with 5.1+, the runtime needs to
// return the host address, as it needs to assume that the host-address is
// device-accessible, as the user has guaranteed it.
// Once the runtime returns the original host-address when the lookup fails, the
// test will need to be updated.
int aa[10][10];
int h[10];
int *ph = &h[0];
@ -24,7 +15,9 @@ struct S {
void f1(int i) {
paa--;
void *original_ph = ph;
void *original_addr_ph3 = &ph[3];
void *original_paa = paa;
void *original_addr_paa102 = &paa[1][0][2];
// (A) No corresponding item, lookup should fail.
@ -34,7 +27,7 @@ struct S {
void *mapped_ptr_ph3 =
omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
mapped_ptr_ph3 != original_addr_ph3, ph == original_ph);
}
// (B) use_device_ptr/map on pointer, and pointee does not exist.
@ -45,7 +38,7 @@ struct S {
void *mapped_ptr_ph3 =
omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr,
mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
mapped_ptr_ph3 != original_addr_ph3, ph == original_ph);
}
// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
@ -78,7 +71,7 @@ struct S {
void *mapped_ptr_paa102 =
omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr,
mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
mapped_ptr_paa102 != original_addr_paa102, paa == original_paa);
}
// (F) use_device_ptr/map on pointer, and pointee does not exist.
@ -89,7 +82,7 @@ struct S {
void *mapped_ptr_paa102 =
omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr,
mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
mapped_ptr_paa102 != original_addr_paa102, paa == original_paa);
}
// (G) map on pointee: base-pointer of map matches use_device_ptr operand.

27
offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
View File

@ -6,15 +6,6 @@
// Test for various cases of use_device_ptr on a reference variable.
// The corresponding data is not previously mapped.
// Note that this tests for the current behavior wherein if a lookup fails,
// the runtime returns nullptr, instead of the original host-address.
// That was compatible with OpenMP 5.0, where it was a user error if
// corresponding storage didn't exist, but with 5.1+, the runtime needs to
// return the host address, as it needs to assume that the host-address is
// device-accessible, as the user has guaranteed it.
// Once the runtime returns the original host-address when the lookup fails, the
// test will need to be updated.
int aa[10][10];
int (*paa_ptee)[10][10] = &aa;
@ -27,32 +18,30 @@ struct S {
void f1(int i) {
paa--;
void *original_ph = ph;
void *original_addr_ph3 = &ph[3];
void *original_paa = paa;
void *original_addr_paa102 = &paa[1][0][2];
// (A) No corresponding item, lookup should fail.
// EXPECTED: A: 1 1 1
// CHECK: A: 1 1 0
// FIXME: ph is not being privatized in the region.
// CHECK: A: 1 1 1
#pragma omp target data use_device_ptr(ph)
{
void *mapped_ptr_ph3 =
omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
mapped_ptr_ph3 != original_addr_ph3, ph == original_ph);
}
// (B) use_device_ptr/map on pointer, and pointee does not exist.
// Lookup should fail.
// EXPECTED: B: 1 1 1
// CHECK: B: 1 1 0
// FIXME: ph is not being privatized in the region.
// CHECK: B: 1 1 1
#pragma omp target data map(ph) use_device_ptr(ph)
{
void *mapped_ptr_ph3 =
omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr,
mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
mapped_ptr_ph3 != original_addr_ph3, ph == original_ph);
}
// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
@ -89,7 +78,7 @@ struct S {
void *mapped_ptr_paa102 =
omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr,
mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
mapped_ptr_paa102 != original_addr_paa102, paa == original_paa);
}
// (F) use_device_ptr/map on pointer, and pointee does not exist.
@ -100,7 +89,7 @@ struct S {
void *mapped_ptr_paa102 =
omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr,
mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
mapped_ptr_paa102 != original_addr_paa102, paa == original_paa);
}
// (G) map on pointee: base-pointer of map matches use_device_ptr operand.

11
offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
View File

@ -7,17 +7,6 @@
// This is necessary because we must assume that the
// pointee is device-accessible, even if it was not
// previously mapped.
//
// OpenMP 5.1, sec 2.14.2, target data construct, p 188, l26-31:
// If a list item that appears in a use_device_ptr clause ... does not point to
// a mapped object, it must contain a valid device address for the target
// device, and the list item references are instead converted to references to a
// local device pointer that refers to this device address.
//
// Note: OpenMP 6.1 will have a way to change the
// fallback behavior: preserve or nullify.
// XFAIL: *
#include <stdio.h>
int x;