[llvm-branch-commits] [clang] [llvm] [Clang][OpenMP] Codegen for `use_device_ptr(fb_nullify)`. (4/4) (PR #173931)
Abhinav Gaba via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Fri Jan 16 11:31:31 PST 2026
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/173931
>From 92109af13f4dd5edd00ed3f00bcb6bc5d118a2a3 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 29 Dec 2025 15:30:46 -0800
Subject: [PATCH 1/4] [Clang][OpenMP] Initial codegen changes for
`use_device_ptr(fb_nullify)`.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 40 ++++++++++++++-----
...vice_ptr_class_member_fallback_nullify.cpp | 4 +-
..._ptr_class_member_ref_fallback_nullify.cpp | 4 +-
...ta_use_device_ptr_var_fallback_nullify.cpp | 4 +-
4 files changed, 33 insertions(+), 19 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 01661ad54ee2f..f0bdf4b6e280f 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7284,6 +7284,7 @@ class MappableExprsHandler {
const ValueDecl *Mapper = nullptr;
const Expr *VarRef = nullptr;
bool ForDeviceAddr = false;
+ bool FbNullify = false;
MapInfo() = default;
MapInfo(
@@ -7293,11 +7294,12 @@ class MappableExprsHandler {
ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
bool ReturnDevicePointer, bool IsImplicit,
const ValueDecl *Mapper = nullptr, const Expr *VarRef = nullptr,
- bool ForDeviceAddr = false)
+ bool ForDeviceAddr = false, bool FbNullify = false)
: Components(Components), MapType(MapType), MapModifiers(MapModifiers),
MotionModifiers(MotionModifiers),
ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit),
- Mapper(Mapper), VarRef(VarRef), ForDeviceAddr(ForDeviceAddr) {}
+ Mapper(Mapper), VarRef(VarRef), ForDeviceAddr(ForDeviceAddr),
+ FbNullify(FbNullify) {}
};
/// The target directive from where the mappable clauses were extracted. It
@@ -8918,7 +8920,8 @@ class MappableExprsHandler {
auto &&UseDeviceDataCombinedInfoGen =
[&UseDeviceDataCombinedInfo](const ValueDecl *VD, llvm::Value *Ptr,
- CodeGenFunction &CGF, bool IsDevAddr) {
+ CodeGenFunction &CGF, bool IsDevAddr,
+ bool FbNullify = false) {
UseDeviceDataCombinedInfo.Exprs.push_back(VD);
UseDeviceDataCombinedInfo.BasePointers.emplace_back(Ptr);
UseDeviceDataCombinedInfo.DevicePtrDecls.emplace_back(VD);
@@ -8932,8 +8935,11 @@ class MappableExprsHandler {
UseDeviceDataCombinedInfo.Pointers.push_back(Ptr);
UseDeviceDataCombinedInfo.Sizes.push_back(
llvm::Constant::getNullValue(CGF.Int64Ty));
- UseDeviceDataCombinedInfo.Types.push_back(
- OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM);
+ OpenMPOffloadMappingFlags Flags =
+ OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM;
+ if (FbNullify)
+ Flags |= OpenMPOffloadMappingFlags::OMP_MAP_FB_NULLIFY;
+ UseDeviceDataCombinedInfo.Types.push_back(Flags);
UseDeviceDataCombinedInfo.Mappers.push_back(nullptr);
};
@@ -8942,7 +8948,8 @@ class MappableExprsHandler {
CodeGenFunction &CGF, const Expr *IE, const ValueDecl *VD,
OMPClauseMappableExprCommon::MappableExprComponentListRef
Components,
- bool IsDevAddr, bool IEIsAttachPtrForDevAddr = false) {
+ bool IsDevAddr, bool IEIsAttachPtrForDevAddr = false,
+ bool FbNullify = false) {
// We didn't find any match in our map information - generate a zero
// size array section.
llvm::Value *Ptr;
@@ -8962,13 +8969,15 @@ class MappableExprsHandler {
// equivalent to
// ... use_device_ptr(p)
UseDeviceDataCombinedInfoGen(VD, Ptr, CGF, /*IsDevAddr=*/IsDevAddr &&
- !TreatDevAddrAsDevPtr);
+ !TreatDevAddrAsDevPtr,
+ FbNullify);
};
auto &&IsMapInfoExist = [&Info, this](CodeGenFunction &CGF,
const ValueDecl *VD, const Expr *IE,
const Expr *DesiredAttachPtrExpr,
- bool IsDevAddr) -> bool {
+ bool IsDevAddr,
+ bool FbNullify = false) -> bool {
// We potentially have map information for this declaration already.
// Look for the first set of components that refer to it. If found,
// return true.
@@ -9000,6 +9009,7 @@ class MappableExprsHandler {
if (IsDevAddr) {
CI->ForDeviceAddr = true;
CI->ReturnDevicePointer = true;
+ CI->FbNullify = FbNullify;
Found = true;
break;
} else {
@@ -9016,6 +9026,7 @@ class MappableExprsHandler {
VD == cast<DeclRefExpr>(AttachPtrExpr)->getDecl())) {
CI->ForDeviceAddr = IsDevAddr;
CI->ReturnDevicePointer = true;
+ CI->FbNullify = FbNullify;
Found = true;
break;
}
@@ -9037,6 +9048,8 @@ class MappableExprsHandler {
const auto *C = dyn_cast<OMPUseDevicePtrClause>(Cl);
if (!C)
continue;
+ bool FbNullify = C->getFallbackModifier() ==
+ OMPC_USE_DEVICE_PTR_FALLBACK_fb_nullify;
for (const auto L : C->component_lists()) {
OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
std::get<1>(L);
@@ -9056,9 +9069,10 @@ class MappableExprsHandler {
Components.front().getAssociatedExpression();
if (IsMapInfoExist(CGF, VD, IE,
/*DesiredAttachPtrExpr=*/UDPOperandExpr,
- /*IsDevAddr=*/false))
+ /*IsDevAddr=*/false, FbNullify))
continue;
- MapInfoGen(CGF, IE, VD, Components, /*IsDevAddr=*/false);
+ MapInfoGen(CGF, IE, VD, Components, /*IsDevAddr=*/false,
+ /*IEIsAttachPtrForDevAddr=*/false, FbNullify);
}
}
@@ -9204,6 +9218,9 @@ class MappableExprsHandler {
: DeviceInfoTy::Pointer;
GroupStructBaseCurInfo.Types[StructBasePointersIdx] |=
OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM;
+ if (L.FbNullify)
+ GroupStructBaseCurInfo.Types[StructBasePointersIdx] |=
+ OpenMPOffloadMappingFlags::OMP_MAP_FB_NULLIFY;
} else {
GroupCurInfo.DevicePtrDecls[CurrentBasePointersIdx] = RelevantVD;
GroupCurInfo.DevicePointers[CurrentBasePointersIdx] =
@@ -9211,6 +9228,9 @@ class MappableExprsHandler {
: DeviceInfoTy::Pointer;
GroupCurInfo.Types[CurrentBasePointersIdx] |=
OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM;
+ if (L.FbNullify)
+ GroupCurInfo.Types[CurrentBasePointersIdx] |=
+ OpenMPOffloadMappingFlags::OMP_MAP_FB_NULLIFY;
}
}
}
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 3094446f8b44d..fca0eeea022b4 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,10 +16,8 @@ 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); // EXPECTED-OFFLOAD-NEXT: (nil)
- // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+ printf("%p\n", a); // OFFLOAD-NEXT: (nil)
// NOOFFLOAD-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 39a987b08a505..65c71738e84ae 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,10 +17,8 @@ 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); // EXPECTED-OFFLOAD-NEXT: (nil)
- // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+ printf("%p\n", b); // OFFLOAD-NEXT: (nil)
// NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
}
};
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 2d4cd11463801..984744cd86bac 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
@@ -13,10 +13,8 @@ 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); // EXPECTED-OFFLOAD-NEXT: (nil)
- // OFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
+ printf("%p\n", xp); // OFFLOAD-NEXT: (nil)
// NOOFFLOAD-NEXT: 0x{{0*}}[[#ADDR]]
}
>From 456d693f3702f7d2914ffde259575effe1e27acf Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 29 Dec 2025 16:11:53 -0800
Subject: [PATCH 2/4] Minor NFC refactor/cleanup.
---
clang/lib/CodeGen/CGOpenMPRuntime.cpp | 73 +++++++++++++--------------
1 file changed, 34 insertions(+), 39 deletions(-)
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index f0bdf4b6e280f..e5e1fbd39d54e 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7284,7 +7284,7 @@ class MappableExprsHandler {
const ValueDecl *Mapper = nullptr;
const Expr *VarRef = nullptr;
bool ForDeviceAddr = false;
- bool FbNullify = false;
+ bool HasUdpFbNullify = false;
MapInfo() = default;
MapInfo(
@@ -7294,12 +7294,12 @@ class MappableExprsHandler {
ArrayRef<OpenMPMotionModifierKind> MotionModifiers,
bool ReturnDevicePointer, bool IsImplicit,
const ValueDecl *Mapper = nullptr, const Expr *VarRef = nullptr,
- bool ForDeviceAddr = false, bool FbNullify = false)
+ bool ForDeviceAddr = false, bool HasUdpFbNullify = false)
: Components(Components), MapType(MapType), MapModifiers(MapModifiers),
MotionModifiers(MotionModifiers),
ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit),
Mapper(Mapper), VarRef(VarRef), ForDeviceAddr(ForDeviceAddr),
- FbNullify(FbNullify) {}
+ HasUdpFbNullify(HasUdpFbNullify) {}
};
/// The target directive from where the mappable clauses were extracted. It
@@ -8921,7 +8921,7 @@ class MappableExprsHandler {
auto &&UseDeviceDataCombinedInfoGen =
[&UseDeviceDataCombinedInfo](const ValueDecl *VD, llvm::Value *Ptr,
CodeGenFunction &CGF, bool IsDevAddr,
- bool FbNullify = false) {
+ bool HasUdpFbNullify = false) {
UseDeviceDataCombinedInfo.Exprs.push_back(VD);
UseDeviceDataCombinedInfo.BasePointers.emplace_back(Ptr);
UseDeviceDataCombinedInfo.DevicePtrDecls.emplace_back(VD);
@@ -8937,7 +8937,7 @@ class MappableExprsHandler {
llvm::Constant::getNullValue(CGF.Int64Ty));
OpenMPOffloadMappingFlags Flags =
OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM;
- if (FbNullify)
+ if (HasUdpFbNullify)
Flags |= OpenMPOffloadMappingFlags::OMP_MAP_FB_NULLIFY;
UseDeviceDataCombinedInfo.Types.push_back(Flags);
UseDeviceDataCombinedInfo.Mappers.push_back(nullptr);
@@ -8949,7 +8949,7 @@ class MappableExprsHandler {
OMPClauseMappableExprCommon::MappableExprComponentListRef
Components,
bool IsDevAddr, bool IEIsAttachPtrForDevAddr = false,
- bool FbNullify = false) {
+ bool HasUdpFbNullify = false) {
// We didn't find any match in our map information - generate a zero
// size array section.
llvm::Value *Ptr;
@@ -8970,14 +8970,13 @@ class MappableExprsHandler {
// ... use_device_ptr(p)
UseDeviceDataCombinedInfoGen(VD, Ptr, CGF, /*IsDevAddr=*/IsDevAddr &&
!TreatDevAddrAsDevPtr,
- FbNullify);
+ HasUdpFbNullify);
};
- auto &&IsMapInfoExist = [&Info, this](CodeGenFunction &CGF,
- const ValueDecl *VD, const Expr *IE,
- const Expr *DesiredAttachPtrExpr,
- bool IsDevAddr,
- bool FbNullify = false) -> bool {
+ auto &&IsMapInfoExist =
+ [&Info, this](CodeGenFunction &CGF, const ValueDecl *VD, const Expr *IE,
+ const Expr *DesiredAttachPtrExpr, bool IsDevAddr,
+ bool HasUdpFbNullify = false) -> bool {
// We potentially have map information for this declaration already.
// Look for the first set of components that refer to it. If found,
// return true.
@@ -9009,7 +9008,7 @@ class MappableExprsHandler {
if (IsDevAddr) {
CI->ForDeviceAddr = true;
CI->ReturnDevicePointer = true;
- CI->FbNullify = FbNullify;
+ CI->HasUdpFbNullify = HasUdpFbNullify;
Found = true;
break;
} else {
@@ -9026,7 +9025,7 @@ class MappableExprsHandler {
VD == cast<DeclRefExpr>(AttachPtrExpr)->getDecl())) {
CI->ForDeviceAddr = IsDevAddr;
CI->ReturnDevicePointer = true;
- CI->FbNullify = FbNullify;
+ CI->HasUdpFbNullify = HasUdpFbNullify;
Found = true;
break;
}
@@ -9048,8 +9047,8 @@ class MappableExprsHandler {
const auto *C = dyn_cast<OMPUseDevicePtrClause>(Cl);
if (!C)
continue;
- bool FbNullify = C->getFallbackModifier() ==
- OMPC_USE_DEVICE_PTR_FALLBACK_fb_nullify;
+ bool HasUdpFbNullify =
+ C->getFallbackModifier() == OMPC_USE_DEVICE_PTR_FALLBACK_fb_nullify;
for (const auto L : C->component_lists()) {
OMPClauseMappableExprCommon::MappableExprComponentListRef Components =
std::get<1>(L);
@@ -9069,10 +9068,10 @@ class MappableExprsHandler {
Components.front().getAssociatedExpression();
if (IsMapInfoExist(CGF, VD, IE,
/*DesiredAttachPtrExpr=*/UDPOperandExpr,
- /*IsDevAddr=*/false, FbNullify))
+ /*IsDevAddr=*/false, HasUdpFbNullify))
continue;
MapInfoGen(CGF, IE, VD, Components, /*IsDevAddr=*/false,
- /*IEIsAttachPtrForDevAddr=*/false, FbNullify);
+ /*IEIsAttachPtrForDevAddr=*/false, HasUdpFbNullify);
}
}
@@ -9209,29 +9208,25 @@ class MappableExprsHandler {
// multiple values are added to any of the lists, the first value
// added is being modified by the assignments below (not the last
// value added).
- if (StructBasePointersIdx <
- GroupStructBaseCurInfo.BasePointers.size()) {
- GroupStructBaseCurInfo.DevicePtrDecls[StructBasePointersIdx] =
- RelevantVD;
- GroupStructBaseCurInfo.DevicePointers[StructBasePointersIdx] =
- L.ForDeviceAddr ? DeviceInfoTy::Address
- : DeviceInfoTy::Pointer;
- GroupStructBaseCurInfo.Types[StructBasePointersIdx] |=
- OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM;
- if (L.FbNullify)
- GroupStructBaseCurInfo.Types[StructBasePointersIdx] |=
- OpenMPOffloadMappingFlags::OMP_MAP_FB_NULLIFY;
- } else {
- GroupCurInfo.DevicePtrDecls[CurrentBasePointersIdx] = RelevantVD;
- GroupCurInfo.DevicePointers[CurrentBasePointersIdx] =
- L.ForDeviceAddr ? DeviceInfoTy::Address
- : DeviceInfoTy::Pointer;
- GroupCurInfo.Types[CurrentBasePointersIdx] |=
+ auto SetDevicePointerInfo = [&](MapCombinedInfoTy &Info,
+ unsigned Idx) {
+ Info.DevicePtrDecls[Idx] = RelevantVD;
+ Info.DevicePointers[Idx] = L.ForDeviceAddr
+ ? DeviceInfoTy::Address
+ : DeviceInfoTy::Pointer;
+ Info.Types[Idx] |=
OpenMPOffloadMappingFlags::OMP_MAP_RETURN_PARAM;
- if (L.FbNullify)
- GroupCurInfo.Types[CurrentBasePointersIdx] |=
+ if (L.HasUdpFbNullify)
+ Info.Types[Idx] |=
OpenMPOffloadMappingFlags::OMP_MAP_FB_NULLIFY;
- }
+ };
+
+ if (StructBasePointersIdx <
+ GroupStructBaseCurInfo.BasePointers.size())
+ SetDevicePointerInfo(GroupStructBaseCurInfo,
+ StructBasePointersIdx);
+ else
+ SetDevicePointerInfo(GroupCurInfo, CurrentBasePointersIdx);
}
}
>From fc10e3bd09bf408fc803c1cb5f84e4776095bf91 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 5 Jan 2026 14:27:14 -0800
Subject: [PATCH 3/4] Add clang codegen test.
---
...t_data_use_device_ptr_fallback_codegen.cpp | 27 +++++++++++++++++++
1 file changed, 27 insertions(+)
create mode 100644 clang/test/OpenMP/target_data_use_device_ptr_fallback_codegen.cpp
diff --git a/clang/test/OpenMP/target_data_use_device_ptr_fallback_codegen.cpp b/clang/test/OpenMP/target_data_use_device_ptr_fallback_codegen.cpp
new file mode 100644
index 0000000000000..0125eecda80c4
--- /dev/null
+++ b/clang/test/OpenMP/target_data_use_device_ptr_fallback_codegen.cpp
@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=61 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -Wno-openmp-mapping -DFB_NULLIFY=1 | FileCheck %s --check-prefix=NULLIFY
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=61 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -Wno-openmp-mapping -DFB_PRESERVE=1 | FileCheck %s --check-prefix=PRESERVE
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=61 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -Wno-openmp-mapping | FileCheck %s --check-prefix=DEFAULT
+
+// expected-no-diagnostics
+
+// NULLIFY: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x8040]]]
+// PRESERVE: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x40]]]
+// DEFAULT: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x40]]]
+
+#ifndef HEADER
+#define HEADER
+
+void f1(void *);
+void f2(int *p) {
+#if FB_NULLIFY
+#pragma omp target data use_device_ptr(fb_nullify: p)
+#elif FB_PRESERVE
+#pragma omp target data use_device_ptr(fb_preserve: p)
+#else
+#pragma omp target data use_device_ptr(p)
+#endif
+ {
+ f1(p);
+ }
+}
+#endif
>From 3834d26cc84f6b06817589fa84edcdcfac47b795 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
Date: Mon, 5 Jan 2026 14:48:34 -0800
Subject: [PATCH 4/4] Update RST files.
---
clang/docs/OpenMPSupport.rst | 4 ++++
clang/docs/ReleaseNotes.rst | 2 ++
2 files changed, 6 insertions(+)
diff --git a/clang/docs/OpenMPSupport.rst b/clang/docs/OpenMPSupport.rst
index 7941c2e439ed6..bdd840ac7922c 100644
--- a/clang/docs/OpenMPSupport.rst
+++ b/clang/docs/OpenMPSupport.rst
@@ -638,6 +638,10 @@ implementation.
| need_device_ptr modifier for adjust_args clause | :part:`partial` | :none:`unclaimed` | Clang Parsing/Sema: https://github.com/llvm/llvm-project/pull/168905 |
| | | | https://github.com/llvm/llvm-project/pull/169558 |
+-------------------------------------------------------------+---------------------------+---------------------------+--------------------------------------------------------------------------+
+| fallback modifier for use_device_ptr clause | :good:`done` | :none:`unclaimed` | Clang: @abhinavgaba (https://github.com/llvm/llvm-project/pull/170578, |
+| | | | https://github.com/llvm/llvm-project/pull/173931) |
+| | | | RT: @abhinavgaba (https://github.com/llvm/llvm-project/pull/169603) |
++-------------------------------------------------------------+---------------------------+---------------------------+--------------------------------------------------------------------------+
OpenMP Extensions
=================
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index d6a2f9e684044..e09fd5bd9ec49 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -990,6 +990,8 @@ OpenMP Support
with OpenMP >= 61.
- ``use_device_ptr`` and ``use_device_addr`` now preserve the original host
address when lookup fails.
+- Added support for ``use_device_ptr`` clause to accept an optional
+ ``fallback`` modifier (``fb_nullify`` or ``fb_preserve``) with OpenMP >= 61.
Improvements
^^^^^^^^^^^^
More information about the llvm-branch-commits
mailing list