[llvm] [OpenMP] Preserve the original address when `use_device_ptr/addr` lookup fails. (PR #169438)
Abhinav Gaba via llvm-commits
llvm-commits at lists.llvm.org
Mon Nov 24 17:08:06 PST 2025
https://github.com/abhinavgaba created https://github.com/llvm/llvm-project/pull/169438
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, so we could do anything
in that scenario.
>From 9824170fed25e52ee9a32b90e9d36a5385733b38 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 24 Nov 2025 13:43:20 -0800
Subject: [PATCH 1/3] [OpenMP] Preserve the original address by default on
use_device_ptr/addr lookup failure.
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, so we could do anything
in that scenario.
---
offload/libomptarget/omptarget.cpp | 34 +++++++++++++++++--
...get_data_use_device_addr_arrsec_fallback.c | 2 --
...target_data_use_device_addr_var_fallback.c | 2 --
.../target_data_use_device_ptr_var_fallback.c | 11 ------
4 files changed, 31 insertions(+), 18 deletions(-)
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 69725e77bae00..3dcc0144f6cf2 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -675,9 +675,37 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsNewEntry ? "" : " not"));
if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
- uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
- void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
- DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
+ intptr_t Delta = reinterpret_cast<intptr_t>(HstPtrBegin) -
+ reinterpret_cast<intptr_t>(HstPtrBase);
+ void *TgtPtrBase;
+ if (TgtPtrBegin) {
+ // Lookup succeeded, return device pointer adjusted by delta
+ TgtPtrBase = reinterpret_cast<void *>(
+ reinterpret_cast<intptr_t>(TgtPtrBegin) - Delta);
+ DP("Returning device pointer " DPxMOD "\n", DPxPTR(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: Support OpenMP 6.1's "fb_nullify" and set the result to
+ // `null - Delta`.
+ TgtPtrBase = reinterpret_cast<void *>(
+ reinterpret_cast<intptr_t>(HstPtrBegin) - Delta);
+ DP("Returning host pointer " DPxMOD " as fallback (lookup failed).\n",
+ DPxPTR(TgtPtrBase));
+ }
ArgsBase[I] = TgtPtrBase;
}
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
index 4b67a3bc2aa7f..118b664fb6e53 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
@@ -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];
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
index 4495a46b6d204..4b0819ef6a9fe 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
@@ -7,8 +7,6 @@
// list-item is device-accessible, even if it was not
// previously mapped.
-// XFAIL: *
-
#include <stdio.h>
int x;
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
index e8fa3b69e9296..33a363495e24a 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
@@ -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;
>From 8e007d1380a31124a46a67f96599bf89d7f00c3e Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 24 Nov 2025 15:49:31 -0800
Subject: [PATCH 2/3] Update some tests that were relying on the previous
behavior.
---
...ta_use_device_addr_arrsec_not_existing.cpp | 20 ++++---------
...se_device_addr_arrsec_ref_not_existing.cpp | 28 +++++--------------
..._data_use_device_addr_var_not_existing.cpp | 21 ++++----------
...a_use_device_addr_var_ref_not_existing.cpp | 21 ++++----------
.../target_wrong_use_device_addr.c | 5 ++--
...arget_data_use_device_ptr_not_existing.cpp | 19 ++++---------
...t_data_use_device_ptr_ref_not_existing.cpp | 27 ++++++------------
7 files changed, 41 insertions(+), 100 deletions(-)
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
index b9ebde431e7bf..78e6bf7c070a0 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
@@ -8,15 +8,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];
@@ -36,7 +27,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.
@@ -58,7 +49,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.
@@ -80,8 +71,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.
@@ -110,7 +100,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.
@@ -122,7 +112,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.
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
index 0090cdb095366..d981da925acc2 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
@@ -8,15 +8,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;
@@ -37,15 +28,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.
@@ -63,15 +52,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.
@@ -95,8 +82,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.
@@ -125,7 +111,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.
@@ -137,7 +123,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.
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
index 79c6f69edba8e..e855b0dd82744 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
@@ -8,15 +8,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];
@@ -38,7 +29,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.
@@ -58,7 +49,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.
@@ -78,7 +69,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.
@@ -99,7 +90,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.
@@ -119,7 +110,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.
@@ -130,7 +121,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.
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
index 9360db4195041..1a3ed148f288b 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
@@ -8,15 +8,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;
@@ -45,7 +36,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.
@@ -65,7 +56,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.
@@ -85,7 +76,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.
@@ -106,7 +97,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.
@@ -126,7 +117,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.
@@ -137,7 +128,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.
diff --git a/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
index 28ec6857fa1a8..f8c9d7c1fe7df 100644
--- a/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
+++ b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
@@ -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);
}
}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
index fe3cdb56e4baa..7632cefb1ea96 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
@@ -8,15 +8,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];
@@ -26,7 +17,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.
@@ -36,7 +29,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.
@@ -47,7 +40,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.
@@ -80,7 +73,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.
@@ -91,7 +84,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.
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
index 419ab3eb33d4d..7c4e18b6bbafd 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
@@ -8,15 +8,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;
@@ -29,32 +20,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.
@@ -91,7 +80,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.
@@ -102,7 +91,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.
>From ef610f43db5f25e2dc1ed8a0471e838f9e006f18 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 24 Nov 2025 16:46:12 -0800
Subject: [PATCH 3/3] Keep using uint64_t.
---
offload/libomptarget/omptarget.cpp | 14 ++++++++------
1 file changed, 8 insertions(+), 6 deletions(-)
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 3dcc0144f6cf2..287564f53101a 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -675,13 +675,13 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsNewEntry ? "" : " not"));
if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
- intptr_t Delta = reinterpret_cast<intptr_t>(HstPtrBegin) -
- reinterpret_cast<intptr_t>(HstPtrBase);
+ 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<intptr_t>(TgtPtrBegin) - Delta);
+ reinterpret_cast<uintptr_t>(TgtPtrBegin) - Delta);
DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
} else {
// Lookup failed. So we have to decide what to do based on the
@@ -699,10 +699,12 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
// to references to a local device pointer that refers to this device
// address.
//
- // TODO: Support OpenMP 6.1's "fb_nullify" and set the result to
- // `null - Delta`.
+ // 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<intptr_t>(HstPtrBegin) - Delta);
+ reinterpret_cast<uintptr_t>(HstPtrBegin) - Delta);
DP("Returning host pointer " DPxMOD " as fallback (lookup failed).\n",
DPxPTR(TgtPtrBase));
}
More information about the llvm-commits
mailing list