https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/169438
>From 9824170fed25e52ee9a32b90e9d36a5385733b38 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Mon, 24 Nov 2025 13:43:20 -0800 Subject: [PATCH 1/4] [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 <[email protected]> Date: Mon, 24 Nov 2025 15:49:31 -0800 Subject: [PATCH 2/4] 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 <[email protected]> Date: Mon, 24 Nov 2025 16:46:12 -0800 Subject: [PATCH 3/4] 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)); } >From 1d76e35bf0115a698ab51b2be195610881e1db56 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Mon, 24 Nov 2025 17:14:49 -0800 Subject: [PATCH 4/4] Update OpenMPSupport.rst, ReleaseNotes.rst. --- clang/docs/OpenMPSupport.rst | 2 ++ clang/docs/ReleaseNotes.rst | 2 ++ 2 files changed, 4 insertions(+) diff --git a/clang/docs/OpenMPSupport.rst b/clang/docs/OpenMPSupport.rst index f7e6061044c6d..7cebf96cfe026 100644 --- a/clang/docs/OpenMPSupport.rst +++ b/clang/docs/OpenMPSupport.rst @@ -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/169438 | ++------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ | device | iterators in map clause or motion clauses | :none:`unclaimed` | | +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ | device | indirect clause on declare target directive | :part:`In Progress` | | diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst index 51f07256c5d9f..ed22cdb39068f 100644 --- a/clang/docs/ReleaseNotes.rst +++ b/clang/docs/ReleaseNotes.rst @@ -759,6 +759,8 @@ OpenMP Support - Updated parsing and semantic analysis support for ``nowait`` clause to accept optional argument in OpenMP >= 60. - Added support for ``default`` clause on ``target`` directive. +- ``use_device_ptr`` and ``use_device_addr`` now preserve the original host + address when lookup fails. Improvements ^^^^^^^^^^^^ _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
