[clang] [llvm] [OpenMP][Clang] Parsing/Sema support for `use_device_ptr(fb_preserve/fb_nullify)`. (2/4) (PR #170578)
Abhinav Gaba via cfe-commits
cfe-commits at lists.llvm.org
Thu Jan 15 13:33:12 PST 2026
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/170578
>From 9824170fed25e52ee9a32b90e9d36a5385733b38 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 24 Nov 2025 13:43:20 -0800
Subject: [PATCH 01/12] [OpenMP] Preserve the original address by default on
use_device_ptr/addr lookup failure.
As per OpenMP 5.1, we need to assume that when the lookup for
use_device_ptr/addr fails, the incoming pointer was already device
accessible.
Prior to 5.1, a lookup-failure meant a user-error, so we could do anything
in that scenario.
---
offload/libomptarget/omptarget.cpp | 34 +++++++++++++++++--
...get_data_use_device_addr_arrsec_fallback.c | 2 --
...target_data_use_device_addr_var_fallback.c | 2 --
.../target_data_use_device_ptr_var_fallback.c | 11 ------
4 files changed, 31 insertions(+), 18 deletions(-)
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 69725e77bae00..3dcc0144f6cf2 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -675,9 +675,37 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsNewEntry ? "" : " not"));
if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
- uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
- void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
- DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
+ intptr_t Delta = reinterpret_cast<intptr_t>(HstPtrBegin) -
+ reinterpret_cast<intptr_t>(HstPtrBase);
+ void *TgtPtrBase;
+ if (TgtPtrBegin) {
+ // Lookup succeeded, return device pointer adjusted by delta
+ TgtPtrBase = reinterpret_cast<void *>(
+ reinterpret_cast<intptr_t>(TgtPtrBegin) - Delta);
+ DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
+ } else {
+ // Lookup failed. So we have to decide what to do based on the
+ // requested fallback behavior.
+ //
+ // Treat "preserve" as the default fallback behavior, since as per
+ // OpenMP 5.1, for use_device_ptr/addr, when there's no corresponding
+ // device pointer to translate into, it's the user's responsibility to
+ // ensure that the host address is device-accessible.
+ //
+ // OpenMP 5.1, sec 2.14.2, target data construct, p 188, l26-31:
+ // If a list item that appears in a use_device_ptr clause ... does not
+ // point to a mapped object, it must contain a valid device address for
+ // the target device, and the list item references are instead converted
+ // to references to a local device pointer that refers to this device
+ // address.
+ //
+ // TODO: Support OpenMP 6.1's "fb_nullify" and set the result to
+ // `null - Delta`.
+ TgtPtrBase = reinterpret_cast<void *>(
+ reinterpret_cast<intptr_t>(HstPtrBegin) - Delta);
+ DP("Returning host pointer " DPxMOD " as fallback (lookup failed).\n",
+ DPxPTR(TgtPtrBase));
+ }
ArgsBase[I] = TgtPtrBase;
}
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
index 4b67a3bc2aa7f..118b664fb6e53 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c
@@ -7,8 +7,6 @@
// list-item is device-accessible, even if it was not
// previously mapped.
-// XFAIL: *
-
#include <stdio.h>
int h[10];
int *ph = &h[0];
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
index 4495a46b6d204..4b0819ef6a9fe 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c
@@ -7,8 +7,6 @@
// list-item is device-accessible, even if it was not
// previously mapped.
-// XFAIL: *
-
#include <stdio.h>
int x;
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
index e8fa3b69e9296..33a363495e24a 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_var_fallback.c
@@ -7,17 +7,6 @@
// This is necessary because we must assume that the
// pointee is device-accessible, even if it was not
// previously mapped.
-//
-// OpenMP 5.1, sec 2.14.2, target data construct, p 188, l26-31:
-// If a list item that appears in a use_device_ptr clause ... does not point to
-// a mapped object, it must contain a valid device address for the target
-// device, and the list item references are instead converted to references to a
-// local device pointer that refers to this device address.
-//
-// Note: OpenMP 6.1 will have a way to change the
-// fallback behavior: preserve or nullify.
-
-// XFAIL: *
#include <stdio.h>
int x;
>From 8e007d1380a31124a46a67f96599bf89d7f00c3e Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 24 Nov 2025 15:49:31 -0800
Subject: [PATCH 02/12] Update some tests that were relying on the previous
behavior.
---
...ta_use_device_addr_arrsec_not_existing.cpp | 20 ++++---------
...se_device_addr_arrsec_ref_not_existing.cpp | 28 +++++--------------
..._data_use_device_addr_var_not_existing.cpp | 21 ++++----------
...a_use_device_addr_var_ref_not_existing.cpp | 21 ++++----------
.../target_wrong_use_device_addr.c | 5 ++--
...arget_data_use_device_ptr_not_existing.cpp | 19 ++++---------
...t_data_use_device_ptr_ref_not_existing.cpp | 27 ++++++------------
7 files changed, 41 insertions(+), 100 deletions(-)
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
index b9ebde431e7bf..78e6bf7c070a0 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp
@@ -8,15 +8,6 @@
// Test for various cases of use_device_addr on an array-section.
// The corresponding data is not previously mapped.
-// Note that this tests for the current behavior wherein if a lookup fails,
-// the runtime returns nullptr, instead of the original host-address.
-// That was compatible with OpenMP 5.0, where it was a user error if
-// corresponding storage didn't exist, but with 5.1+, the runtime needs to
-// return the host address, as it needs to assume that the host-address is
-// device-accessible, as the user has guaranteed it.
-// Once the runtime returns the original host-address when the lookup fails, the
-// test will need to be updated.
-
int g, h[10];
int *ph = &h[0];
@@ -36,7 +27,7 @@ struct S {
int *mapped_ptr_ph3 =
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (B) use_device_addr/map: different operands, same base-pointer.
@@ -58,7 +49,7 @@ struct S {
int *mapped_ptr_ph3 =
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (D) use_device_addr/map: one of two maps with matching base-pointer.
@@ -80,8 +71,7 @@ struct S {
int **mapped_ptr_paa02 =
(int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
- mapped_ptr_paa02 != original_paa02,
- &paa[0][2] == (int **)nullptr + 2);
+ mapped_ptr_paa02 != original_paa02, &paa[0][2] == original_paa02);
}
// (F) use_device_addr/map: different operands, same base-array.
@@ -110,7 +100,7 @@ struct S {
}
int *original_paa020 = &paa[0][2][0];
- int **original_paa0 = (int **)&paa[0];
+ void *original_paa0 = &paa[0];
// (H) use_device_addr/map: different base-pointers.
// No corresponding storage for use_device_addr opnd, lookup should fail.
@@ -122,7 +112,7 @@ struct S {
int **mapped_ptr_paa0 =
(int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
- mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
+ mapped_ptr_paa0 == nullptr, &paa[0] == original_paa0);
}
// (I) use_device_addr/map: one map with different, one with same base-ptr.
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
index 0090cdb095366..d981da925acc2 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp
@@ -8,15 +8,6 @@
// Test for various cases of use_device_addr on an array-section on a reference.
// The corresponding data is not previously mapped.
-// Note that this tests for the current behavior wherein if a lookup fails,
-// the runtime returns nullptr, instead of the original host-address.
-// That was compatible with OpenMP 5.0, where it was a user error if
-// corresponding storage didn't exist, but with 5.1+, the runtime needs to
-// return the host address, as it needs to assume that the host-address is
-// device-accessible, as the user has guaranteed it.
-// Once the runtime returns the original host-address when the lookup fails, the
-// test will need to be updated.
-
int g_ptee;
int &g = g_ptee;
@@ -37,15 +28,13 @@ struct S {
int **original_paa02 = &paa[0][2];
// (A) No corresponding map, lookup should fail.
-// EXPECTED: A: 1 1 1
-// CHECK: A: 1 1 0
-// FIXME: ph is not being privatized in the region.
+// CHECK: A: 1 1 1
#pragma omp target data use_device_addr(ph[3 : 4])
{
int *mapped_ptr_ph3 =
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (B) use_device_addr/map: different operands, same base-pointer.
@@ -63,15 +52,13 @@ struct S {
// (C) use_device_addr/map: different base-pointers.
// No corresponding storage, lookup should fail.
-// EXPECTED: C: 1 1 1
-// CHECK: C: 1 1 0
-// FIXME: ph is not being privatized in the region.
+// CHECK: C: 1 1 1
#pragma omp target data map(ph) use_device_addr(ph[3 : 4])
{
int *mapped_ptr_ph3 =
(int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
}
// (D) use_device_addr/map: one of two maps with matching base-pointer.
@@ -95,8 +82,7 @@ struct S {
int **mapped_ptr_paa02 =
(int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
- mapped_ptr_paa02 != original_paa02,
- &paa[0][2] == (int **)nullptr + 2);
+ mapped_ptr_paa02 != original_paa02, &paa[0][2] == original_paa02);
}
// (F) use_device_addr/map: different operands, same base-array.
@@ -125,7 +111,7 @@ struct S {
}
int *original_paa020 = &paa[0][2][0];
- int **original_paa0 = (int **)&paa[0];
+ void *original_paa0 = &paa[0];
// (H) use_device_addr/map: different base-pointers.
// No corresponding storage for use_device_addr opnd, lookup should fail.
@@ -137,7 +123,7 @@ struct S {
int **mapped_ptr_paa0 =
(int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
- mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
+ mapped_ptr_paa0 == nullptr, &paa[0] == original_paa0);
}
// (I) use_device_addr/map: one map with different, one with same base-ptr.
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
index 79c6f69edba8e..e855b0dd82744 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp
@@ -8,15 +8,6 @@
// Test for various cases of use_device_addr on a variable (not a section).
// The corresponding data is not previously mapped.
-// Note that this tests for the current behavior wherein if a lookup fails,
-// the runtime returns nullptr, instead of the original host-address.
-// That was compatible with OpenMP 5.0, where it was a user error if
-// corresponding storage didn't exist, but with 5.1+, the runtime needs to
-// return the host address, as it needs to assume that the host-address is
-// device-accessible, as the user has guaranteed it.
-// Once the runtime returns the original host-address when the lookup fails, the
-// test will need to be updated.
-
int g, h[10];
int *ph = &h[0];
@@ -38,7 +29,7 @@ struct S {
void *mapped_ptr_g =
omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_g == nullptr,
- mapped_ptr_g != original_addr_g, (void *)&g == nullptr);
+ mapped_ptr_g != original_addr_g, &g == original_addr_g);
}
// (B) Lookup should succeed.
@@ -58,7 +49,7 @@ struct S {
void *mapped_ptr_h =
omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
printf("C: %d %d %d\n", mapped_ptr_h == nullptr,
- mapped_ptr_h != original_addr_h, (void *)&h == nullptr);
+ mapped_ptr_h != original_addr_h, &h == original_addr_h);
}
// (D) Lookup should succeed.
@@ -78,7 +69,7 @@ struct S {
void *mapped_ptr_ph =
omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_ph == nullptr,
- mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (F) Lookup should succeed.
@@ -99,7 +90,7 @@ struct S {
void *mapped_ptr_ph =
omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
printf("G: %d %d %d\n", mapped_ptr_ph == nullptr,
- mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (H) Maps both pointee and pointer. Lookup for pointer should succeed.
@@ -119,7 +110,7 @@ struct S {
void *mapped_ptr_paa =
omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
printf("I: %d %d %d\n", mapped_ptr_paa == nullptr,
- mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (J) Maps pointee only, but use_device_addr operand is pointer.
@@ -130,7 +121,7 @@ struct S {
void *mapped_ptr_paa =
omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
printf("J: %d %d %d\n", mapped_ptr_paa == nullptr,
- mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (K) Lookup should succeed.
diff --git a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
index 9360db4195041..1a3ed148f288b 100644
--- a/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
+++ b/offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp
@@ -8,15 +8,6 @@
// Test for various cases of use_device_addr on a reference variable.
// The corresponding data is not previously mapped.
-// Note that this tests for the current behavior wherein if a lookup fails,
-// the runtime returns nullptr, instead of the original host-address.
-// That was compatible with OpenMP 5.0, where it was a user error if
-// corresponding storage didn't exist, but with 5.1+, the runtime needs to
-// return the host address, as it needs to assume that the host-address is
-// device-accessible, as the user has guaranteed it.
-// Once the runtime returns the original host-address when the lookup fails, the
-// test will need to be updated.
-
int g_ptee;
int &g = g_ptee;
@@ -45,7 +36,7 @@ struct S {
void *mapped_ptr_g =
omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_g == nullptr,
- mapped_ptr_g != original_addr_g, (void *)&g == nullptr);
+ mapped_ptr_g != original_addr_g, &g == original_addr_g);
}
// (B) Lookup should succeed.
@@ -65,7 +56,7 @@ struct S {
void *mapped_ptr_h =
omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
printf("C: %d %d %d\n", mapped_ptr_h == nullptr,
- mapped_ptr_h != original_addr_h, (void *)&h == nullptr);
+ mapped_ptr_h != original_addr_h, &h == original_addr_h);
}
// (D) Lookup should succeed.
@@ -85,7 +76,7 @@ struct S {
void *mapped_ptr_ph =
omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_ph == nullptr,
- mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (F) Lookup should succeed.
@@ -106,7 +97,7 @@ struct S {
void *mapped_ptr_ph =
omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
printf("G: %d %d %d\n", mapped_ptr_ph == nullptr,
- mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
}
// (H) Maps both pointee and pointer. Lookup for pointer should succeed.
@@ -126,7 +117,7 @@ struct S {
void *mapped_ptr_paa =
omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
printf("I: %d %d %d\n", mapped_ptr_paa == nullptr,
- mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (J) Maps pointee only, but use_device_addr operand is pointer.
@@ -137,7 +128,7 @@ struct S {
void *mapped_ptr_paa =
omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
printf("J: %d %d %d\n", mapped_ptr_paa == nullptr,
- mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
}
// (K) Lookup should succeed.
diff --git a/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
index 28ec6857fa1a8..f8c9d7c1fe7df 100644
--- a/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
+++ b/offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c
@@ -1,5 +1,5 @@
// RUN: %libomptarget-compile-generic -fopenmp-version=51 -g
-// RUN: env LIBOMPTARGET_INFO=64 %libomptarget-run-fail-generic 2>&1 \
+// RUN: env LIBOMPTARGET_INFO=64 %libomptarget-run-generic 2>&1 \
// RUN: | %fcheck-generic
// FIXME: Fails due to optimized debugging in 'ptxas'
@@ -20,7 +20,8 @@ int main() {
// counterpart
#pragma omp target data use_device_addr(x)
{
- // CHECK-NOT: device addr=0x[[#%x,HOST_ADDR:]]
+ // Even when the lookup fails, x should retain its host address.
+ // CHECK: device addr=0x[[#HOST_ADDR]]
fprintf(stderr, "device addr=%p\n", x);
}
}
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
index fe3cdb56e4baa..7632cefb1ea96 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_not_existing.cpp
@@ -8,15 +8,6 @@
// Test for various cases of use_device_ptr on a variable.
// The corresponding data is not previously mapped.
-// Note that this tests for the current behavior wherein if a lookup fails,
-// the runtime returns nullptr, instead of the original host-address.
-// That was compatible with OpenMP 5.0, where it was a user error if
-// corresponding storage didn't exist, but with 5.1+, the runtime needs to
-// return the host address, as it needs to assume that the host-address is
-// device-accessible, as the user has guaranteed it.
-// Once the runtime returns the original host-address when the lookup fails, the
-// test will need to be updated.
-
int aa[10][10];
int h[10];
int *ph = &h[0];
@@ -26,7 +17,9 @@ struct S {
void f1(int i) {
paa--;
+ void *original_ph = ph;
void *original_addr_ph3 = &ph[3];
+ void *original_paa = paa;
void *original_addr_paa102 = &paa[1][0][2];
// (A) No corresponding item, lookup should fail.
@@ -36,7 +29,7 @@ struct S {
void *mapped_ptr_ph3 =
omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ mapped_ptr_ph3 != original_addr_ph3, ph == original_ph);
}
// (B) use_device_ptr/map on pointer, and pointee does not exist.
@@ -47,7 +40,7 @@ struct S {
void *mapped_ptr_ph3 =
omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ mapped_ptr_ph3 != original_addr_ph3, ph == original_ph);
}
// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
@@ -80,7 +73,7 @@ struct S {
void *mapped_ptr_paa102 =
omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr,
- mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ mapped_ptr_paa102 != original_addr_paa102, paa == original_paa);
}
// (F) use_device_ptr/map on pointer, and pointee does not exist.
@@ -91,7 +84,7 @@ struct S {
void *mapped_ptr_paa102 =
omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr,
- mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ mapped_ptr_paa102 != original_addr_paa102, paa == original_paa);
}
// (G) map on pointee: base-pointer of map matches use_device_ptr operand.
diff --git a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
index 419ab3eb33d4d..7c4e18b6bbafd 100644
--- a/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
+++ b/offload/test/mapping/use_device_ptr/target_data_use_device_ptr_ref_not_existing.cpp
@@ -8,15 +8,6 @@
// Test for various cases of use_device_ptr on a reference variable.
// The corresponding data is not previously mapped.
-// Note that this tests for the current behavior wherein if a lookup fails,
-// the runtime returns nullptr, instead of the original host-address.
-// That was compatible with OpenMP 5.0, where it was a user error if
-// corresponding storage didn't exist, but with 5.1+, the runtime needs to
-// return the host address, as it needs to assume that the host-address is
-// device-accessible, as the user has guaranteed it.
-// Once the runtime returns the original host-address when the lookup fails, the
-// test will need to be updated.
-
int aa[10][10];
int (*paa_ptee)[10][10] = &aa;
@@ -29,32 +20,30 @@ struct S {
void f1(int i) {
paa--;
+ void *original_ph = ph;
void *original_addr_ph3 = &ph[3];
+ void *original_paa = paa;
void *original_addr_paa102 = &paa[1][0][2];
// (A) No corresponding item, lookup should fail.
-// EXPECTED: A: 1 1 1
-// CHECK: A: 1 1 0
-// FIXME: ph is not being privatized in the region.
+// CHECK: A: 1 1 1
#pragma omp target data use_device_ptr(ph)
{
void *mapped_ptr_ph3 =
omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ mapped_ptr_ph3 != original_addr_ph3, ph == original_ph);
}
// (B) use_device_ptr/map on pointer, and pointee does not exist.
// Lookup should fail.
-// EXPECTED: B: 1 1 1
-// CHECK: B: 1 1 0
-// FIXME: ph is not being privatized in the region.
+// CHECK: B: 1 1 1
#pragma omp target data map(ph) use_device_ptr(ph)
{
void *mapped_ptr_ph3 =
omp_get_mapped_ptr(original_addr_ph3, omp_get_default_device());
printf("B: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_addr_ph3, ph == nullptr);
+ mapped_ptr_ph3 != original_addr_ph3, ph == original_ph);
}
// (C) map on pointee: base-pointer of map matches use_device_ptr operand.
@@ -91,7 +80,7 @@ struct S {
void *mapped_ptr_paa102 =
omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
printf("E: %d %d %d\n", mapped_ptr_paa102 == nullptr,
- mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ mapped_ptr_paa102 != original_addr_paa102, paa == original_paa);
}
// (F) use_device_ptr/map on pointer, and pointee does not exist.
@@ -102,7 +91,7 @@ struct S {
void *mapped_ptr_paa102 =
omp_get_mapped_ptr(original_addr_paa102, omp_get_default_device());
printf("F: %d %d %d\n", mapped_ptr_paa102 == nullptr,
- mapped_ptr_paa102 != original_addr_paa102, paa == nullptr);
+ mapped_ptr_paa102 != original_addr_paa102, paa == original_paa);
}
// (G) map on pointee: base-pointer of map matches use_device_ptr operand.
>From ef610f43db5f25e2dc1ed8a0471e838f9e006f18 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 24 Nov 2025 16:46:12 -0800
Subject: [PATCH 03/12] 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 <abhinav.gaba at intel.com>
Date: Mon, 24 Nov 2025 17:14:49 -0800
Subject: [PATCH 04/12] 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 <abhinav.gaba at intel.com>
Date: Tue, 25 Nov 2025 15:31:22 -0800
Subject: [PATCH 05/12] [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 <abhinav.gaba at intel.com>
Date: Mon, 1 Dec 2025 13:22:40 -0800
Subject: [PATCH 06/12] 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 <abhinav.gaba at intel.com>
Date: Tue, 23 Dec 2025 15:38:48 -0800
Subject: [PATCH 07/12] 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 <abhinav.gaba at intel.com>
Date: Tue, 6 Jan 2026 13:59:31 -0800
Subject: [PATCH 08/12] 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 <abhinav.gaba at intel.com>
Date: Mon, 1 Dec 2025 14:20:15 -0800
Subject: [PATCH 09/12] [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 <abhinav.gaba at intel.com>
Date: Wed, 3 Dec 2025 16:06:54 -0800
Subject: [PATCH 10/12] 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 <abhinav.gaba at intel.com>
Date: Wed, 3 Dec 2025 16:32:12 -0800
Subject: [PATCH 11/12] 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 25dfbc6a1030792eae6a6c782f9149e3abd4432e Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Thu, 15 Jan 2026 13:32:18 -0800
Subject: [PATCH 12/12] Add error-checking test for alternating modifiers and
list-items in same clause.
---
.../OpenMP/target_data_use_device_ptr_fallback_messages.cpp | 4 ++++
1 file changed, 4 insertions(+)
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 fff2dcf15e29e..ae6fda3ce0939 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
@@ -22,6 +22,10 @@ void f1(int x, int *p, int *q) {
#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'}}
{}
+ // Interspersed modifiers/list-items: should fail in both versions
+#pragma omp target data map(x) use_device_ptr(fb_nullify: p, fb_preserve: q) // 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'}}
{}
More information about the cfe-commits
mailing list