https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/173930
>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 01/15] [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 02/15] 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 03/15] 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 04/15] 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 ^^^^^^^^^^^^ >From 3fd3927df233e887d8a2e9133c0c22ab07c66487 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Tue, 25 Nov 2025 15:31:22 -0800 Subject: [PATCH 05/15] [OpenMP][Offload] Add `FB_NULLIFY` map-type for `use_device_ptr(fb_nullify)`. This PR adds a new map-type bit to control the fallback behavior when when a pointer lookup fails. For now, this is only meaningful with `RETURN_PARAM`, and can be used for `need_device_ptr` (for which the default is to use `nullptr` as the result when lookup fails), and OpenMP 6.1's `use_device_ptr(fb_nullify)`. Eventually, this can be extended to work with assumed-size maps on `target` constructs, to control what the argument should be set to when lookup fails (the OpenMP spec does not have a way to control that yet). --- .../llvm/Frontend/OpenMP/OMPConstants.h | 4 ++++ offload/include/omptarget.h | 4 ++++ offload/libomptarget/omptarget.cpp | 22 ++++++++++++------- 3 files changed, 22 insertions(+), 8 deletions(-) diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h b/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h index 58fd8a490c04a..d2a1b5209ecba 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPConstants.h @@ -252,6 +252,10 @@ enum class OpenMPOffloadMappingFlags : uint64_t { // Attach pointer and pointee, after processing all other maps. // Applicable to map-entering directives. Does not change ref-count. OMP_MAP_ATTACH = 0x4000, + // When a lookup fails, fall back to using null as the translated pointer, + // instead of preserving the original pointer's value. Currently only + // useful in conjunction with RETURN_PARAM. + OMP_MAP_FB_NULLIFY = 0x8000, /// Signal that the runtime library should use args as an array of /// descriptor_dim pointers and use args_size as dims. Used when we have /// non-contiguous list items in target update directive diff --git a/offload/include/omptarget.h b/offload/include/omptarget.h index fbb4a06accf84..44e19a5290c48 100644 --- a/offload/include/omptarget.h +++ b/offload/include/omptarget.h @@ -80,6 +80,10 @@ enum tgt_map_type { // Attach pointer and pointee, after processing all other maps. // Applicable to map-entering directives. Does not change ref-count. OMP_TGT_MAPTYPE_ATTACH = 0x4000, + // When a lookup fails, fall back to using null as the translated pointer, + // instead of preserving the original pointer's value. Currently only + // useful in conjunction with RETURN_PARAM. + OMP_TGT_MAPTYPE_FB_NULLIFY = 0x8000, // descriptor for non-contiguous target-update OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000, // member of struct, member given by [16 MSBs] - 1 diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index 287564f53101a..d2376a527c1da 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -699,14 +699,20 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, // 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); - DP("Returning host pointer " DPxMOD " as fallback (lookup failed).\n", - DPxPTR(TgtPtrBase)); + // OpenMP 6.1's `fb_nullify` fallback behavior: when the FB_NULLIFY bit + // is set by the compiler, e.g. for `use/need_device_ptr(fb_nullify)`), + // return `nullptr - Delta` when lookup fails. + if (ArgTypes[I] & OMP_TGT_MAPTYPE_FB_NULLIFY) { + TgtPtrBase = reinterpret_cast<void *>( + reinterpret_cast<uintptr_t>(nullptr) - Delta); + DP("Returning offsetted null pointer " DPxMOD " as fallback (lookup failed)\n", + DPxPTR(TgtPtrBase)); + } else { + TgtPtrBase = reinterpret_cast<void *>( + reinterpret_cast<uintptr_t>(HstPtrBegin) - Delta); + DP("Returning host pointer " DPxMOD " as fallback (lookup failed)\n", + DPxPTR(TgtPtrBase)); + } } ArgsBase[I] = TgtPtrBase; } >From aa6a1b74f2000a16bafdec6db481babdd3e752bf Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Mon, 1 Dec 2025 13:22:40 -0800 Subject: [PATCH 06/15] Clang-format --- offload/libomptarget/omptarget.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index d2376a527c1da..669f8d0b1a85c 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -705,7 +705,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, if (ArgTypes[I] & OMP_TGT_MAPTYPE_FB_NULLIFY) { TgtPtrBase = reinterpret_cast<void *>( reinterpret_cast<uintptr_t>(nullptr) - Delta); - DP("Returning offsetted null pointer " DPxMOD " as fallback (lookup failed)\n", + DP("Returning offsetted null pointer " DPxMOD + " as fallback (lookup failed)\n", DPxPTR(TgtPtrBase)); } else { TgtPtrBase = reinterpret_cast<void *>( >From e716fa8120309d5b8780cd2ea77a84231e351c88 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Tue, 23 Dec 2025 15:38:48 -0800 Subject: [PATCH 07/15] Fix minor typo. --- offload/libomptarget/omptarget.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp index b6b03e692ad26..960c5bc17df96 100644 --- a/offload/libomptarget/omptarget.cpp +++ b/offload/libomptarget/omptarget.cpp @@ -713,9 +713,8 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum, if (ArgTypes[I] & OMP_TGT_MAPTYPE_FB_NULLIFY) { TgtPtrBase = reinterpret_cast<void *>( reinterpret_cast<uintptr_t>(nullptr) - Delta); - ODBG(ODT_MAPPING) << "Returning offsetted null pointer " - - << TgtPtrBase << " as fallback (lookup failed)"; + ODBG(ODT_Mapping) << "Returning offsetted null pointer " << TgtPtrBase + << " as fallback (lookup failed)"; } else { TgtPtrBase = reinterpret_cast<void *>( reinterpret_cast<uintptr_t>(HstPtrBegin) - Delta); >From 7ae746ff1cfd0230f080c43aa36c3996398f00cc Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Tue, 6 Jan 2026 13:59:31 -0800 Subject: [PATCH 08/15] Update PR number in OpenMPSupport RST. --- clang/docs/OpenMPSupport.rst | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/docs/OpenMPSupport.rst b/clang/docs/OpenMPSupport.rst index 273e7ea601c8c..7941c2e439ed6 100644 --- a/clang/docs/OpenMPSupport.rst +++ b/clang/docs/OpenMPSupport.rst @@ -266,7 +266,7 @@ 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 | 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 | +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ >From 009a8ae6c1939f86c7ac3b6ddcb011a403b84e45 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Mon, 1 Dec 2025 14:20:15 -0800 Subject: [PATCH 09/15] [OpenMP][Clang] Parsing/Sema support for `use_device_ptr(fb_preserve/fb_nullify)`. Depends on #169603. This is the `use_device_ptr` counterpart of #168905. With OpenMP 6.1, a `fallback` modifier can be specified on the `use_device_ptr` clause to control the behavior when a pointer lookup fails, i.e. there is no device pointer to translate into. The default is `fb_preserve` (i.e. retain the original pointer), while `fb_nullify` means: use `nullptr` as the translated pointer. --- clang/include/clang/AST/OpenMPClause.h | 39 ++++++++++++++++++++--- clang/include/clang/Basic/OpenMPKinds.def | 8 +++++ clang/include/clang/Basic/OpenMPKinds.h | 8 +++++ clang/include/clang/Sema/SemaOpenMP.h | 8 +++-- clang/lib/AST/OpenMPClause.cpp | 17 ++++++++-- clang/lib/Basic/OpenMPKinds.cpp | 22 +++++++++++-- clang/lib/Parse/ParseOpenMP.cpp | 18 +++++++++++ clang/lib/Sema/SemaOpenMP.cpp | 14 +++++--- clang/lib/Sema/TreeTransform.h | 12 ++++--- clang/lib/Serialization/ASTReader.cpp | 2 ++ clang/lib/Serialization/ASTWriter.cpp | 2 ++ 11 files changed, 130 insertions(+), 20 deletions(-) diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h index 6525e64ff102f..0847839221ea0 100644 --- a/clang/include/clang/AST/OpenMPClause.h +++ b/clang/include/clang/AST/OpenMPClause.h @@ -7989,6 +7989,13 @@ class OMPUseDevicePtrClause final friend OMPVarListClause; friend TrailingObjects; + /// Fallback modifier for the clause. + OpenMPUseDevicePtrFallbackModifier FallbackModifier = + OMPC_USE_DEVICE_PTR_FALLBACK_unknown; + + /// Location of the fallback modifier. + SourceLocation FallbackModifierLoc; + /// Build clause with number of variables \a NumVars. /// /// \param Locs Locations needed to build a mappable clause. It includes 1) @@ -7999,10 +8006,14 @@ class OMPUseDevicePtrClause final /// NumUniqueDeclarations: number of unique base declarations in this clause; /// 3) NumComponentLists: number of component lists in this clause; and 4) /// NumComponents: total number of expression components in the clause. + /// \param FallbackModifier The fallback modifier for the clause. + /// \param FallbackModifierLoc Location of the fallback modifier. explicit OMPUseDevicePtrClause(const OMPVarListLocTy &Locs, - const OMPMappableExprListSizeTy &Sizes) - : OMPMappableExprListClause(llvm::omp::OMPC_use_device_ptr, Locs, Sizes) { - } + const OMPMappableExprListSizeTy &Sizes, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc) + : OMPMappableExprListClause(llvm::omp::OMPC_use_device_ptr, Locs, Sizes), + FallbackModifier(FallbackModifier), FallbackModifierLoc(FallbackModifierLoc) {} /// Build an empty clause. /// @@ -8055,6 +8066,14 @@ class OMPUseDevicePtrClause final return {getPrivateCopies().end(), varlist_size()}; } + /// Set the fallback modifier for the clause. + void setFallbackModifier(OpenMPUseDevicePtrFallbackModifier M) { + FallbackModifier = M; + } + + /// Set the location of the fallback modifier. + void setFallbackModifierLoc(SourceLocation Loc) { FallbackModifierLoc = Loc; } + public: /// Creates clause with a list of variables \a Vars. /// @@ -8067,11 +8086,15 @@ class OMPUseDevicePtrClause final /// \param Inits Expressions referring to private copy initializers. /// \param Declarations Declarations used in the clause. /// \param ComponentLists Component lists used in the clause. + /// \param FallbackModifier The fallback modifier for the clause. + /// \param FallbackModifierLoc Location of the fallback modifier. static OMPUseDevicePtrClause * Create(const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars, ArrayRef<Expr *> PrivateVars, ArrayRef<Expr *> Inits, ArrayRef<ValueDecl *> Declarations, - MappableExprComponentListsRef ComponentLists); + MappableExprComponentListsRef ComponentLists, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc); /// Creates an empty clause with the place for \a NumVars variables. /// @@ -8084,6 +8107,14 @@ class OMPUseDevicePtrClause final static OMPUseDevicePtrClause * CreateEmpty(const ASTContext &C, const OMPMappableExprListSizeTy &Sizes); + /// Get the fallback modifier for the clause. + OpenMPUseDevicePtrFallbackModifier getFallbackModifier() const { + return FallbackModifier; + } + + /// Get the location of the fallback modifier. + SourceLocation getFallbackModifierLoc() const { return FallbackModifierLoc; } + using private_copies_iterator = MutableArrayRef<Expr *>::iterator; using private_copies_const_iterator = ArrayRef<const Expr *>::iterator; using private_copies_range = llvm::iterator_range<private_copies_iterator>; diff --git a/clang/include/clang/Basic/OpenMPKinds.def b/clang/include/clang/Basic/OpenMPKinds.def index ceac89d3aba6d..e61ee0ddc08da 100644 --- a/clang/include/clang/Basic/OpenMPKinds.def +++ b/clang/include/clang/Basic/OpenMPKinds.def @@ -110,6 +110,9 @@ #ifndef OPENMP_NEED_DEVICE_PTR_KIND #define OPENMP_NEED_DEVICE_PTR_KIND(Name) #endif +#ifndef OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER +#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name) +#endif // Static attributes for 'schedule' clause. OPENMP_SCHEDULE_KIND(static) @@ -282,6 +285,10 @@ OPENMP_THREADSET_KIND(omp_team) OPENMP_NEED_DEVICE_PTR_KIND(fb_nullify) OPENMP_NEED_DEVICE_PTR_KIND(fb_preserve) +// OpenMP 6.1 modifiers for 'use_device_ptr' clause. +OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(fb_nullify) +OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(fb_preserve) + #undef OPENMP_NUMTASKS_MODIFIER #undef OPENMP_NUMTHREADS_MODIFIER #undef OPENMP_DYN_GROUPPRIVATE_MODIFIER @@ -315,3 +322,4 @@ OPENMP_NEED_DEVICE_PTR_KIND(fb_preserve) #undef OPENMP_ALLOCATE_MODIFIER #undef OPENMP_THREADSET_KIND #undef OPENMP_NEED_DEVICE_PTR_KIND +#undef OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER diff --git a/clang/include/clang/Basic/OpenMPKinds.h b/clang/include/clang/Basic/OpenMPKinds.h index 3b088b3efd998..4e83bfcd0128b 100644 --- a/clang/include/clang/Basic/OpenMPKinds.h +++ b/clang/include/clang/Basic/OpenMPKinds.h @@ -218,6 +218,14 @@ enum OpenMPNeedDevicePtrModifier { OMPC_NEED_DEVICE_PTR_unknown, }; +/// OpenMP 6.1 use_device_ptr fallback modifier +enum OpenMPUseDevicePtrFallbackModifier { +#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name) \ + OMPC_USE_DEVICE_PTR_FALLBACK_##Name, +#include "clang/Basic/OpenMPKinds.def" + OMPC_USE_DEVICE_PTR_FALLBACK_unknown, +}; + /// OpenMP bindings for the 'bind' clause. enum OpenMPBindClauseKind { #define OPENMP_BIND_KIND(Name) OMPC_BIND_##Name, diff --git a/clang/include/clang/Sema/SemaOpenMP.h b/clang/include/clang/Sema/SemaOpenMP.h index 2d05b4423140b..e4eb3345534a4 100644 --- a/clang/include/clang/Sema/SemaOpenMP.h +++ b/clang/include/clang/Sema/SemaOpenMP.h @@ -1176,6 +1176,9 @@ class SemaOpenMP : public SemaBase { int OriginalSharingModifier = 0; // Default is shared int NeedDevicePtrModifier = 0; SourceLocation NeedDevicePtrModifierLoc; + int UseDevicePtrFallbackModifier = + OMPC_USE_DEVICE_PTR_FALLBACK_unknown; ///< Fallback modifier for use_device_ptr clause. + SourceLocation UseDevicePtrFallbackModifierLoc; SmallVector<OpenMPMapModifierKind, NumberOfOMPMapClauseModifiers> MapTypeModifiers; SmallVector<SourceLocation, NumberOfOMPMapClauseModifiers> @@ -1364,8 +1367,9 @@ class SemaOpenMP : public SemaBase { ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers = {}); /// Called on well-formed 'use_device_ptr' clause. - OMPClause *ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList, - const OMPVarListLocTy &Locs); + OMPClause *ActOnOpenMPUseDevicePtrClause( + ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, SourceLocation FallbackModifierLoc); /// Called on well-formed 'use_device_addr' clause. OMPClause *ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs); diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp index 2183d77de8fa7..5a6a958595671 100644 --- a/clang/lib/AST/OpenMPClause.cpp +++ b/clang/lib/AST/OpenMPClause.cpp @@ -1441,7 +1441,9 @@ OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create( const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars, ArrayRef<Expr *> PrivateVars, ArrayRef<Expr *> Inits, ArrayRef<ValueDecl *> Declarations, - MappableExprComponentListsRef ComponentLists) { + MappableExprComponentListsRef ComponentLists, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc) { OMPMappableExprListSizeTy Sizes; Sizes.NumVars = Vars.size(); Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations); @@ -1465,7 +1467,8 @@ OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create( Sizes.NumUniqueDeclarations + Sizes.NumComponentLists, Sizes.NumComponents)); - OMPUseDevicePtrClause *Clause = new (Mem) OMPUseDevicePtrClause(Locs, Sizes); + OMPUseDevicePtrClause *Clause = new (Mem) + OMPUseDevicePtrClause(Locs, Sizes, FallbackModifier, FallbackModifierLoc); Clause->setVarRefs(Vars); Clause->setPrivateCopies(PrivateVars); @@ -2753,7 +2756,15 @@ void OMPClausePrinter::VisitOMPDefaultmapClause(OMPDefaultmapClause *Node) { void OMPClausePrinter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *Node) { if (!Node->varlist_empty()) { OS << "use_device_ptr"; - VisitOMPClauseList(Node, '('); + if (Node->getFallbackModifier() != OMPC_USE_DEVICE_PTR_FALLBACK_unknown) { + OS << "(" + << getOpenMPSimpleClauseTypeName(OMPC_use_device_ptr, + Node->getFallbackModifier()) + << ":"; + VisitOMPClauseList(Node, ' '); + } else { + VisitOMPClauseList(Node, '('); + } OS << ")"; } } diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp index 03485b7e81abc..7ba2c89638c05 100644 --- a/clang/lib/Basic/OpenMPKinds.cpp +++ b/clang/lib/Basic/OpenMPKinds.cpp @@ -238,6 +238,16 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str, return OMPC_NUMTHREADS_unknown; return Type; } + case OMPC_use_device_ptr: { + unsigned Type = llvm::StringSwitch<unsigned>(Str) +#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name) \ + .Case(#Name, OMPC_USE_DEVICE_PTR_FALLBACK_##Name) +#include "clang/Basic/OpenMPKinds.def" + .Default(OMPC_USE_DEVICE_PTR_FALLBACK_unknown); + if (LangOpts.OpenMP < 61) + return OMPC_USE_DEVICE_PTR_FALLBACK_unknown; + return Type; + } case OMPC_unknown: case OMPC_threadprivate: case OMPC_groupprivate: @@ -280,7 +290,6 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str, case OMPC_nogroup: case OMPC_hint: case OMPC_uniform: - case OMPC_use_device_ptr: case OMPC_use_device_addr: case OMPC_is_device_ptr: case OMPC_has_device_addr: @@ -608,6 +617,16 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind, #include "clang/Basic/OpenMPKinds.def" } llvm_unreachable("Invalid OpenMP 'threadset' clause modifier"); + case OMPC_use_device_ptr: + switch (Type) { + case OMPC_USE_DEVICE_PTR_FALLBACK_unknown: + return "unknown"; +#define OPENMP_USE_DEVICE_PTR_FALLBACK_MODIFIER(Name) \ + case OMPC_USE_DEVICE_PTR_FALLBACK_##Name: \ + return #Name; +#include "clang/Basic/OpenMPKinds.def" + } + llvm_unreachable("Invalid OpenMP 'use_device_ptr' clause modifier"); case OMPC_unknown: case OMPC_threadprivate: case OMPC_groupprivate: @@ -650,7 +669,6 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind, case OMPC_nogroup: case OMPC_hint: case OMPC_uniform: - case OMPC_use_device_ptr: case OMPC_use_device_addr: case OMPC_is_device_ptr: case OMPC_has_device_addr: diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp index 01fd05961f876..3619cd03ec0fd 100644 --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -5055,6 +5055,24 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind, ExpectAndConsume(tok::colon, diag::warn_pragma_expected_colon, "adjust-op"); } + } else if (Kind == OMPC_use_device_ptr && getLangOpts().OpenMP >= 61) { + // Handle optional fallback modifier for use_device_ptr clause. + // use_device_ptr([fb_preserve | fb_nullify :] list) + // Default is fb_preserve. + if (Tok.is(tok::identifier)) { + auto FallbackModifier = static_cast<OpenMPUseDevicePtrFallbackModifier>( + getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok), getLangOpts())); + if (FallbackModifier != OMPC_USE_DEVICE_PTR_FALLBACK_unknown) { + Data.UseDevicePtrFallbackModifier = FallbackModifier; + Data.UseDevicePtrFallbackModifierLoc = Tok.getLocation(); + ConsumeToken(); + if (Tok.is(tok::colon)) { + Data.ColonLoc = ConsumeToken(); + } else { + Diag(Tok, diag::err_modifier_expected_colon) << "fallback"; + } + } + } } bool IsComma = diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 2a1337be13b99..29c809888a168 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -18757,7 +18757,11 @@ OMPClause *SemaOpenMP::ActOnOpenMPVarListClause(OpenMPClauseKind Kind, VarList, Locs); break; case OMPC_use_device_ptr: - Res = ActOnOpenMPUseDevicePtrClause(VarList, Locs); + Res = ActOnOpenMPUseDevicePtrClause( + VarList, Locs, + static_cast<OpenMPUseDevicePtrFallbackModifier>( + Data.UseDevicePtrFallbackModifier), + Data.UseDevicePtrFallbackModifierLoc); break; case OMPC_use_device_addr: Res = ActOnOpenMPUseDeviceAddrClause(VarList, Locs); @@ -24574,9 +24578,9 @@ OMPClause *SemaOpenMP::ActOnOpenMPFromClause( MapperId); } -OMPClause * -SemaOpenMP::ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList, - const OMPVarListLocTy &Locs) { +OMPClause *SemaOpenMP::ActOnOpenMPUseDevicePtrClause( + ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, SourceLocation FallbackModifierLoc) { MappableVarListInfo MVLI(VarList); SmallVector<Expr *, 8> PrivateCopies; SmallVector<Expr *, 8> Inits; @@ -24657,7 +24661,7 @@ SemaOpenMP::ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList, return OMPUseDevicePtrClause::Create( getASTContext(), Locs, MVLI.ProcessedVarList, PrivateCopies, Inits, - MVLI.VarBaseDeclarations, MVLI.VarComponents); + MVLI.VarBaseDeclarations, MVLI.VarComponents, FallbackModifier, FallbackModifierLoc); } OMPClause * diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index bc923c80b7132..25ed92fd3f44b 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -2258,9 +2258,12 @@ class TreeTransform { /// /// By default, performs semantic analysis to build the new OpenMP clause. /// Subclasses may override this routine to provide different behavior. - OMPClause *RebuildOMPUseDevicePtrClause(ArrayRef<Expr *> VarList, - const OMPVarListLocTy &Locs) { - return getSema().OpenMP().ActOnOpenMPUseDevicePtrClause(VarList, Locs); + OMPClause *RebuildOMPUseDevicePtrClause( + ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc) { + return getSema().OpenMP().ActOnOpenMPUseDevicePtrClause( + VarList, Locs, FallbackModifier, FallbackModifierLoc); } /// Build a new OpenMP 'use_device_addr' clause. @@ -11624,7 +11627,8 @@ OMPClause *TreeTransform<Derived>::TransformOMPUseDevicePtrClause( Vars.push_back(EVar.get()); } OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc()); - return getDerived().RebuildOMPUseDevicePtrClause(Vars, Locs); + return getDerived().RebuildOMPUseDevicePtrClause( + Vars, Locs, C->getFallbackModifier(), C->getFallbackModifierLoc()); } template <typename Derived> diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 66cf484bb5cb6..b6b1a4d280b16 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -12542,6 +12542,8 @@ void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) { void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) { C->setLParenLoc(Record.readSourceLocation()); + C->setFallbackModifier(Record.readEnum<OpenMPUseDevicePtrFallbackModifier>()); + C->setFallbackModifierLoc(Record.readSourceLocation()); auto NumVars = C->varlist_size(); auto UniqueDecls = C->getUniqueDeclarationsNum(); auto TotalLists = C->getTotalComponentListNum(); diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp index 39104da10d0b7..d66dc7b2adffd 100644 --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -8535,6 +8535,8 @@ void OMPClauseWriter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) { Record.push_back(C->getTotalComponentListNum()); Record.push_back(C->getTotalComponentsNum()); Record.AddSourceLocation(C->getLParenLoc()); + Record.writeEnum(C->getFallbackModifier()); + Record.AddSourceLocation(C->getFallbackModifierLoc()); for (auto *E : C->varlist()) Record.AddStmt(E); for (auto *VE : C->private_copies()) >From 0e260ad20c41ed834f7615d9509f6e46add1cb4e Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Wed, 3 Dec 2025 16:06:54 -0800 Subject: [PATCH 10/15] Add tests. --- ...data_use_device_ptr_fallback_ast_print.cpp | 36 +++++++++++++++++++ ..._data_use_device_ptr_fallback_messages.cpp | 28 +++++++++++++++ 2 files changed, 64 insertions(+) create mode 100644 clang/test/OpenMP/target_data_use_device_ptr_fallback_ast_print.cpp create mode 100644 clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp diff --git a/clang/test/OpenMP/target_data_use_device_ptr_fallback_ast_print.cpp b/clang/test/OpenMP/target_data_use_device_ptr_fallback_ast_print.cpp new file mode 100644 index 0000000000000..060f64f6e86a8 --- /dev/null +++ b/clang/test/OpenMP/target_data_use_device_ptr_fallback_ast_print.cpp @@ -0,0 +1,36 @@ +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=61 -ast-print %s | FileCheck %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=61 -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=61 -std=c++11 -include-pch %t -verify %s -ast-print | FileCheck %s + +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +// CHECK-LABEL:void f1(int *p, int *q) +void f1(int *p, int *q) { + +// CHECK: #pragma omp target data use_device_ptr(fb_preserve: p) +#pragma omp target data use_device_ptr(fb_preserve: p) + {} + +// CHECK: #pragma omp target data use_device_ptr(fb_nullify: p) +#pragma omp target data use_device_ptr(fb_nullify: p) + {} + +// Without any fallback modifier +// CHECK: #pragma omp target data use_device_ptr(p) +#pragma omp target data use_device_ptr(p) + {} + +// Multiple variables with fb_preserve +// CHECK: #pragma omp target data use_device_ptr(fb_preserve: p,q) +#pragma omp target data use_device_ptr(fb_preserve: p, q) + {} + +// Multiple variables with fb_nullify +// CHECK: #pragma omp target data use_device_ptr(fb_nullify: p,q) +#pragma omp target data use_device_ptr(fb_nullify: p, q) + {} +} +#endif diff --git a/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp b/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp new file mode 100644 index 0000000000000..7a22e95e7fee6 --- /dev/null +++ b/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp @@ -0,0 +1,28 @@ +// RUN: %clang_cc1 -std=c++11 -fopenmp -fopenmp-version=60 -verify=omp60,expected -ferror-limit 200 %s +// RUN: %clang_cc1 -std=c++11 -fopenmp -fopenmp-version=61 -verify=omp61,expected -ferror-limit 200 %s + +void f1(int x, int *p, int *q) { + + // Test that fallback modifier is only recognized in OpenMP 6.1+ +#pragma omp target data map(x) use_device_ptr(fb_preserve: p) // omp60-error {{use of undeclared identifier 'fb_preserve'}} + {} + +#pragma omp target data map(x) use_device_ptr(fb_nullify: p) // omp60-error {{use of undeclared identifier 'fb_nullify'}} + {} + + // Without modifier (should work in both versions) +#pragma omp target data map(x) use_device_ptr(p) + {} + + // Unknown modifier: should fail in both versions +#pragma omp target data map(x) use_device_ptr(fb_abc: p) // expected-error {{use of undeclared identifier 'fb_abc'}} + {} + + // Multiple modifiers: should fail in both versions +#pragma omp target data map(x) use_device_ptr(fb_nullify, fb_preserve: p, q) // omp61-error {{missing ':' after fallback modifier}} omp61-error {{expected expression}} omp61-error {{use of undeclared identifier 'fb_preserve'}} omp60-error {{use of undeclared identifier 'fb_nullify'}} omp60-error {{use of undeclared identifier 'fb_preserve'}} + {} + + // Test missing colon after modifier in OpenMP 6.1 - should error +#pragma omp target data map(x) use_device_ptr(fb_preserve p) // omp61-error {{missing ':' after fallback modifier}} omp60-error {{use of undeclared identifier 'fb_preserve'}} + {} +} \ No newline at end of file >From 1986d1fd37eba52fb726875bc44373ca70f5b316 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Wed, 3 Dec 2025 16:32:12 -0800 Subject: [PATCH 11/15] Minor formatting changes. --- clang/include/clang/AST/OpenMPClause.h | 11 ++++++----- clang/include/clang/Sema/SemaOpenMP.h | 6 ++++-- clang/lib/Parse/ParseOpenMP.cpp | 5 ++--- clang/lib/Sema/SemaOpenMP.cpp | 6 ++++-- .../target_data_use_device_ptr_fallback_messages.cpp | 2 +- 5 files changed, 17 insertions(+), 13 deletions(-) diff --git a/clang/include/clang/AST/OpenMPClause.h b/clang/include/clang/AST/OpenMPClause.h index 0847839221ea0..21a4cfb519f5a 100644 --- a/clang/include/clang/AST/OpenMPClause.h +++ b/clang/include/clang/AST/OpenMPClause.h @@ -8008,12 +8008,13 @@ class OMPUseDevicePtrClause final /// NumComponents: total number of expression components in the clause. /// \param FallbackModifier The fallback modifier for the clause. /// \param FallbackModifierLoc Location of the fallback modifier. - explicit OMPUseDevicePtrClause(const OMPVarListLocTy &Locs, - const OMPMappableExprListSizeTy &Sizes, - OpenMPUseDevicePtrFallbackModifier FallbackModifier, - SourceLocation FallbackModifierLoc) + explicit OMPUseDevicePtrClause( + const OMPVarListLocTy &Locs, const OMPMappableExprListSizeTy &Sizes, + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc) : OMPMappableExprListClause(llvm::omp::OMPC_use_device_ptr, Locs, Sizes), - FallbackModifier(FallbackModifier), FallbackModifierLoc(FallbackModifierLoc) {} + FallbackModifier(FallbackModifier), + FallbackModifierLoc(FallbackModifierLoc) {} /// Build an empty clause. /// diff --git a/clang/include/clang/Sema/SemaOpenMP.h b/clang/include/clang/Sema/SemaOpenMP.h index e4eb3345534a4..1d4ea0f1cf3b0 100644 --- a/clang/include/clang/Sema/SemaOpenMP.h +++ b/clang/include/clang/Sema/SemaOpenMP.h @@ -1177,7 +1177,8 @@ class SemaOpenMP : public SemaBase { int NeedDevicePtrModifier = 0; SourceLocation NeedDevicePtrModifierLoc; int UseDevicePtrFallbackModifier = - OMPC_USE_DEVICE_PTR_FALLBACK_unknown; ///< Fallback modifier for use_device_ptr clause. + OMPC_USE_DEVICE_PTR_FALLBACK_unknown; ///< Fallback modifier for + ///< use_device_ptr clause. SourceLocation UseDevicePtrFallbackModifierLoc; SmallVector<OpenMPMapModifierKind, NumberOfOMPMapClauseModifiers> MapTypeModifiers; @@ -1369,7 +1370,8 @@ class SemaOpenMP : public SemaBase { /// Called on well-formed 'use_device_ptr' clause. OMPClause *ActOnOpenMPUseDevicePtrClause( ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, - OpenMPUseDevicePtrFallbackModifier FallbackModifier, SourceLocation FallbackModifierLoc); + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc); /// Called on well-formed 'use_device_addr' clause. OMPClause *ActOnOpenMPUseDeviceAddrClause(ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs); diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp index 3619cd03ec0fd..9e181bd5bfb1d 100644 --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -5066,11 +5066,10 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind, Data.UseDevicePtrFallbackModifier = FallbackModifier; Data.UseDevicePtrFallbackModifierLoc = Tok.getLocation(); ConsumeToken(); - if (Tok.is(tok::colon)) { + if (Tok.is(tok::colon)) Data.ColonLoc = ConsumeToken(); - } else { + else Diag(Tok, diag::err_modifier_expected_colon) << "fallback"; - } } } } diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 29c809888a168..f6c6b493d819b 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -24580,7 +24580,8 @@ OMPClause *SemaOpenMP::ActOnOpenMPFromClause( OMPClause *SemaOpenMP::ActOnOpenMPUseDevicePtrClause( ArrayRef<Expr *> VarList, const OMPVarListLocTy &Locs, - OpenMPUseDevicePtrFallbackModifier FallbackModifier, SourceLocation FallbackModifierLoc) { + OpenMPUseDevicePtrFallbackModifier FallbackModifier, + SourceLocation FallbackModifierLoc) { MappableVarListInfo MVLI(VarList); SmallVector<Expr *, 8> PrivateCopies; SmallVector<Expr *, 8> Inits; @@ -24661,7 +24662,8 @@ OMPClause *SemaOpenMP::ActOnOpenMPUseDevicePtrClause( return OMPUseDevicePtrClause::Create( getASTContext(), Locs, MVLI.ProcessedVarList, PrivateCopies, Inits, - MVLI.VarBaseDeclarations, MVLI.VarComponents, FallbackModifier, FallbackModifierLoc); + MVLI.VarBaseDeclarations, MVLI.VarComponents, FallbackModifier, + FallbackModifierLoc); } OMPClause * diff --git a/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp b/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp index 7a22e95e7fee6..fff2dcf15e29e 100644 --- a/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp +++ b/clang/test/OpenMP/target_data_use_device_ptr_fallback_messages.cpp @@ -25,4 +25,4 @@ void f1(int x, int *p, int *q) { // Test missing colon after modifier in OpenMP 6.1 - should error #pragma omp target data map(x) use_device_ptr(fb_preserve p) // omp61-error {{missing ':' after fallback modifier}} omp60-error {{use of undeclared identifier 'fb_preserve'}} {} -} \ No newline at end of file +} >From 7dbda3107a06748a12ce044746d093cdfcc41e30 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Fri, 5 Dec 2025 15:10:27 -0800 Subject: [PATCH 12/15] [NFC][OpenMP][Offload] Add tests for `use_device_ptr(fb_preserve/nullify)`. The fallback modifiers are currently part of OpenMP 6.1. The tests mostly fail for now. The associated libomptarget and clang parsing/sema changes are in #169438, #169603 and #170578, with clang codegen to follow. --- ...a_use_device_ptr_class_member_fallback.cpp | 34 ++++++++++++++++++ ...vice_ptr_class_member_fallback_nullify.cpp | 30 ++++++++++++++++ ...ice_ptr_class_member_fallback_preserve.cpp | 30 ++++++++++++++++ ...e_device_ptr_class_member_ref_fallback.cpp | 35 +++++++++++++++++++ ..._ptr_class_member_ref_fallback_nullify.cpp | 31 ++++++++++++++++ ...ptr_class_member_ref_fallback_preserve.cpp | 31 ++++++++++++++++ ...rget_data_use_device_ptr_var_fallback.cpp} | 6 +++- ...ta_use_device_ptr_var_fallback_nullify.cpp | 23 ++++++++++++ ...a_use_device_ptr_var_fallback_preserve.cpp | 23 ++++++++++++ ...t_data_use_device_ptr_var_ref_fallback.cpp | 26 ++++++++++++++ ...se_device_ptr_var_ref_fallback_nullify.cpp | 25 +++++++++++++ ...e_device_ptr_var_ref_fallback_preserve.cpp | 24 +++++++++++++ 12 files changed, 317 insertions(+), 1 deletion(-) create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp rename offload/test/mapping/use_device_ptr/{target_data_use_device_ptr_var_fallback.c => target_data_use_device_ptr_var_fallback.cpp} (69%) create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp create mode 100644 offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp new file mode 100644 index 0000000000000..2dd33732a7d2f --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp @@ -0,0 +1,34 @@ +// RUN: %libomptarget-compilexx-generic +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value by +// default. +// +// This is necessary because we must assume that the +// pointee is device-accessible, even if it was not +// previously mapped. + +// XFAIL: * + +#include <stdio.h> + +int x = 0; + +struct ST { + int *a = &x; + + void f1() { + printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(a) + printf("%p\n", a); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f1(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp new file mode 100644 index 0000000000000..61f3367f537ee --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp @@ -0,0 +1,30 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer is set to null because of fb_nullify. + +// XFAIL: * + +#include <stdio.h> + +int x = 0; + +struct ST { + int *a = &x; + + void f1() { + printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_nullify : a) + printf("%p\n", a); // OFFLOAD-NEXT: (nil) + // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f1(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp new file mode 100644 index 0000000000000..b7af1f39cc3bf --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp @@ -0,0 +1,30 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value +// because of fb_preserve. + +// XFAIL: * + +#include <stdio.h> + +int x = 0; + +struct ST { + int *a = &x; + + void f1() { + printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_preserve : a) + printf("%p\n", a); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f1(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp new file mode 100644 index 0000000000000..45f89d0ee92cc --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp @@ -0,0 +1,35 @@ +// RUN: %libomptarget-compilexx-generic +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value by +// default. +// +// This is necessary because we must assume that the +// pointee is device-accessible, even if it was not +// previously mapped. + +// XFAIL: * + +#include <stdio.h> + +int x = 0; +int *y = &x; + +struct ST { + int *&b = y; + + void f2() { + printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(b) + printf("%p\n", b); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f2(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp new file mode 100644 index 0000000000000..39f39a974577c --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp @@ -0,0 +1,31 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer is set to null because of fb_nullify. + +// XFAIL: * + +#include <stdio.h> + +int x = 0; +int *y = &x; + +struct ST { + int *&b = y; + + void f2() { + printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_nullify : b) + printf("%p\n", b); // OFFLOAD-NEXT: (nil) + // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f2(); +} diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp new file mode 100644 index 0000000000000..ad861ab12001e --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp @@ -0,0 +1,31 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value +// because of fb_preserve. + +// XFAIL: * + +#include <stdio.h> + +int x = 0; +int *y = &x; + +struct ST { + int *&b = y; + + void f2() { + printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_preserve : b) + printf("%p\n", b); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] + } +}; + +int main() { + ST s; + s.f2(); +} 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.cpp similarity index 69% rename from offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c rename to offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.cpp index 33a363495e24a..5be209a8c5b0d 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.cpp @@ -1,4 +1,8 @@ -// RUN: %libomptarget-compilexx-run-and-check-generic +// RUN: %libomptarget-compilexx-generic +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic // Test that when a use_device_ptr lookup fails, the // privatized pointer retains its original value by diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp new file mode 100644 index 0000000000000..cb11f52645dd8 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp @@ -0,0 +1,23 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer is set to null because of fb_nullify. + +// XFAIL: * + +#include <stdio.h> +int x; +int *xp = &x; + +void f1() { + printf("%p\n", xp); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_nullify : xp) + printf("%p\n", xp); // OFFLOAD-NEXT: (nil) + // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f1(); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp new file mode 100644 index 0000000000000..31ce803fc1ed0 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp @@ -0,0 +1,23 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value +// because of fb_preserve. + +// XFAIL: * + +#include <stdio.h> +int x; +int *xp = &x; + +void f1() { + printf("%p\n", xp); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_preserve : xp) + printf("%p\n", xp); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f1(); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp new file mode 100644 index 0000000000000..1060ed9cdbc70 --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback.cpp @@ -0,0 +1,26 @@ +// RUN: %libomptarget-compilexx-generic +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value by +// default. +// +// This is necessary because we must assume that the +// pointee is device-accessible, even if it was not +// previously mapped. + +#include <stdio.h> +int x; +int *xp = &x; +int *&xpr = xp; + +void f2() { + printf("%p\n", xpr); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(xpr) + printf("%p\n", xpr); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f2(); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp new file mode 100644 index 0000000000000..230ffda4fad9a --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp @@ -0,0 +1,25 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,OFFLOAD +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic -check-prefixes=CHECK,NOOFFLOAD + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer is set to null because of fb_nullify. + +// XFAIL: * + +#include <stdio.h> +int x; +int *xp = &x; +int *&xpr = xp; + +void f2() { + printf("%p\n", xpr); // CHECK: 0x[[#%x,ADDR:]] + // FIXME: We won't get "nil" until we start privatizing xpr. +#pragma omp target data use_device_ptr(fb_nullify : xpr) + printf("%p\n", xpr); // OFFLOAD-NEXT: (nil) + // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f2(); } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp new file mode 100644 index 0000000000000..443739814ed0d --- /dev/null +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp @@ -0,0 +1,24 @@ +// RUN: %libomptarget-compilexx-generic -fopenmp-version=61 +// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic +// RUN: env OMP_TARGET_OFFLOAD=disabled %libomptarget-run-generic 2>&1 \ +// RUN: | %fcheck-generic + +// Test that when a use_device_ptr lookup fails, the +// privatized pointer retains its original value +// because of fb_preserve. + +// XFAIL: * + +#include <stdio.h> +int x; +int *xp = &x; +int *&xpr = xp; + +void f2() { + printf("%p\n", xpr); // CHECK: 0x[[#%x,ADDR:]] +#pragma omp target data use_device_ptr(fb_preserve : xpr) + printf("%p\n", xpr); // CHECK-NEXT: 0x{{0*}}[[#ADDR]] +} + +int main() { f2(); } >From 75d82a2d78cf959b9ab4d5833d5ee043e33e1bfb Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Mon, 29 Dec 2025 15:03:49 -0800 Subject: [PATCH 13/15] Update test checks to incorporate Parsing/Sema changes. --- .../target_data_use_device_ptr_class_member_fallback.cpp | 2 -- ...et_data_use_device_ptr_class_member_fallback_nullify.cpp | 6 +++--- ...t_data_use_device_ptr_class_member_fallback_preserve.cpp | 2 -- ...target_data_use_device_ptr_class_member_ref_fallback.cpp | 2 -- ...ata_use_device_ptr_class_member_ref_fallback_nullify.cpp | 6 +++--- ...ta_use_device_ptr_class_member_ref_fallback_preserve.cpp | 2 -- .../target_data_use_device_ptr_var_fallback_nullify.cpp | 6 +++--- .../target_data_use_device_ptr_var_fallback_preserve.cpp | 2 -- .../target_data_use_device_ptr_var_ref_fallback_nullify.cpp | 5 ++--- ...target_data_use_device_ptr_var_ref_fallback_preserve.cpp | 2 -- 10 files changed, 11 insertions(+), 24 deletions(-) diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp index 2dd33732a7d2f..5c232d5db02e0 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback.cpp @@ -12,8 +12,6 @@ // pointee is device-accessible, even if it was not // previously mapped. -// XFAIL: * - #include <stdio.h> int x = 0; diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp index 61f3367f537ee..e85f51736f4f7 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp @@ -7,8 +7,6 @@ // Test that when a use_device_ptr lookup fails, the // privatized pointer is set to null because of fb_nullify. -// XFAIL: * - #include <stdio.h> int x = 0; @@ -18,8 +16,10 @@ struct ST { void f1() { printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]] + // FIXME: Update this with codegen changes for fb_nullify #pragma omp target data use_device_ptr(fb_nullify : a) - printf("%p\n", a); // OFFLOAD-NEXT: (nil) + printf("%p\n", a); // EXPECTED-OFFLOAD-NEXT: (nil) + // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] } }; diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp index b7af1f39cc3bf..51944c561dc2d 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_preserve.cpp @@ -8,8 +8,6 @@ // privatized pointer retains its original value // because of fb_preserve. -// XFAIL: * - #include <stdio.h> int x = 0; diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp index 45f89d0ee92cc..59a8facdf2f5c 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback.cpp @@ -12,8 +12,6 @@ // pointee is device-accessible, even if it was not // previously mapped. -// XFAIL: * - #include <stdio.h> int x = 0; diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp index 39f39a974577c..00e6372a0f588 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp @@ -7,8 +7,6 @@ // Test that when a use_device_ptr lookup fails, the // privatized pointer is set to null because of fb_nullify. -// XFAIL: * - #include <stdio.h> int x = 0; @@ -19,8 +17,10 @@ struct ST { void f2() { printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]] + // FIXME: Update this with codegen changes for fb_nullify #pragma omp target data use_device_ptr(fb_nullify : b) - printf("%p\n", b); // OFFLOAD-NEXT: (nil) + printf("%p\n", b); // EXPECTED-OFFLOAD-NEXT: (nil) + // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] } }; diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp index ad861ab12001e..beeb7526e1625 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_preserve.cpp @@ -8,8 +8,6 @@ // privatized pointer retains its original value // because of fb_preserve. -// XFAIL: * - #include <stdio.h> int x = 0; diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp index cb11f52645dd8..2d4cd11463801 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_nullify.cpp @@ -7,16 +7,16 @@ // Test that when a use_device_ptr lookup fails, the // privatized pointer is set to null because of fb_nullify. -// XFAIL: * - #include <stdio.h> int x; int *xp = &x; void f1() { printf("%p\n", xp); // CHECK: 0x[[#%x,ADDR:]] + // FIXME: Update this with codegen changes for fb_nullify #pragma omp target data use_device_ptr(fb_nullify : xp) - printf("%p\n", xp); // OFFLOAD-NEXT: (nil) + printf("%p\n", xp); // EXPECTED-OFFLOAD-NEXT: (nil) + // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp index 31ce803fc1ed0..197704f14f86a 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback_preserve.cpp @@ -8,8 +8,6 @@ // privatized pointer retains its original value // because of fb_preserve. -// XFAIL: * - #include <stdio.h> int x; int *xp = &x; diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp index 230ffda4fad9a..7fa76dd69e7c0 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_nullify.cpp @@ -7,8 +7,6 @@ // Test that when a use_device_ptr lookup fails, the // privatized pointer is set to null because of fb_nullify. -// XFAIL: * - #include <stdio.h> int x; int *xp = &x; @@ -18,7 +16,8 @@ void f2() { printf("%p\n", xpr); // CHECK: 0x[[#%x,ADDR:]] // FIXME: We won't get "nil" until we start privatizing xpr. #pragma omp target data use_device_ptr(fb_nullify : xpr) - printf("%p\n", xpr); // OFFLOAD-NEXT: (nil) + printf("%p\n", xpr); // EXPECTED-OFFLOAD-NEXT: (nil) + // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] // NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] } diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp index 443739814ed0d..e7f8bd48ec4fe 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_ref_fallback_preserve.cpp @@ -8,8 +8,6 @@ // privatized pointer retains its original value // because of fb_preserve. -// XFAIL: * - #include <stdio.h> int x; int *xp = &x; >From 660e4edf398a3840e823f9eb93abe4ff1b35954d Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Mon, 29 Dec 2025 15:21:04 -0800 Subject: [PATCH 14/15] Clang-format fixes. --- ...target_data_use_device_ptr_class_member_fallback_nullify.cpp | 2 +- ...et_data_use_device_ptr_class_member_ref_fallback_nullify.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp index e85f51736f4f7..9745276294078 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp @@ -16,7 +16,7 @@ struct ST { void f1() { printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]] - // FIXME: Update this with codegen changes for fb_nullify + // FIXME: Update this with codegen changes for fb_nullify #pragma omp target data use_device_ptr(fb_nullify : a) printf("%p\n", a); // EXPECTED-OFFLOAD-NEXT: (nil) // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp index 00e6372a0f588..76610a95af512 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp @@ -17,7 +17,7 @@ struct ST { void f2() { printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]] - // FIXME: Update this with codegen changes for fb_nullify + // FIXME: Update this with codegen changes for fb_nullify #pragma omp target data use_device_ptr(fb_nullify : b) printf("%p\n", b); // EXPECTED-OFFLOAD-NEXT: (nil) // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] >From 80310d75873a37273438529a97d68893f26c8d67 Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Tue, 30 Dec 2025 13:37:37 -0800 Subject: [PATCH 15/15] Clang-format changed its mind. --- ...target_data_use_device_ptr_class_member_fallback_nullify.cpp | 2 +- ...et_data_use_device_ptr_class_member_ref_fallback_nullify.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp index 9745276294078..3094446f8b44d 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_fallback_nullify.cpp @@ -16,7 +16,7 @@ struct ST { void f1() { printf("%p\n", a); // CHECK: 0x[[#%x,ADDR:]] - // FIXME: Update this with codegen changes for fb_nullify + // FIXME: Update this with codegen changes for fb_nullify #pragma omp target data use_device_ptr(fb_nullify : a) printf("%p\n", a); // EXPECTED-OFFLOAD-NEXT: (nil) // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp index 76610a95af512..39a987b08a505 100644 --- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp +++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_class_member_ref_fallback_nullify.cpp @@ -17,7 +17,7 @@ struct ST { void f2() { printf("%p\n", b); // CHECK: 0x[[#%x,ADDR:]] - // FIXME: Update this with codegen changes for fb_nullify + // FIXME: Update this with codegen changes for fb_nullify #pragma omp target data use_device_ptr(fb_nullify : b) printf("%p\n", b); // EXPECTED-OFFLOAD-NEXT: (nil) // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]] _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
