[llvm-branch-commits] [clang] [llvm] [WiP][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
Mon Dec 29 16:17:28 PST 2025
https://github.com/abhinavgaba created https://github.com/llvm/llvm-project/pull/173931
TBD: CG LIT tests, rst file updates.
>From f3309d108f4c52c103f2238075f83d8c86b5685d 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/2] [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 b8ee701c482bb..91d3983f6c284 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7162,6 +7162,7 @@ class MappableExprsHandler {
const ValueDecl *Mapper = nullptr;
const Expr *VarRef = nullptr;
bool ForDeviceAddr = false;
+ bool FbNullify = false;
MapInfo() = default;
MapInfo(
@@ -7171,11 +7172,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
@@ -8796,7 +8798,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);
@@ -8810,8 +8813,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);
};
@@ -8820,7 +8826,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;
@@ -8840,13 +8847,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.
@@ -8878,6 +8887,7 @@ class MappableExprsHandler {
if (IsDevAddr) {
CI->ForDeviceAddr = true;
CI->ReturnDevicePointer = true;
+ CI->FbNullify = FbNullify;
Found = true;
break;
} else {
@@ -8894,6 +8904,7 @@ class MappableExprsHandler {
VD == cast<DeclRefExpr>(AttachPtrExpr)->getDecl())) {
CI->ForDeviceAddr = IsDevAddr;
CI->ReturnDevicePointer = true;
+ CI->FbNullify = FbNullify;
Found = true;
break;
}
@@ -8915,6 +8926,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);
@@ -8934,9 +8947,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);
}
}
@@ -9082,6 +9096,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] =
@@ -9089,6 +9106,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 9745276294078..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 76610a95af512..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 7574ae5968fb1f9f26b8152272e7ad065f1a2e7a 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/2] 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 91d3983f6c284..e16ef36a98374 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7162,7 +7162,7 @@ class MappableExprsHandler {
const ValueDecl *Mapper = nullptr;
const Expr *VarRef = nullptr;
bool ForDeviceAddr = false;
- bool FbNullify = false;
+ bool HasUdpFbNullify = false;
MapInfo() = default;
MapInfo(
@@ -7172,12 +7172,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
@@ -8799,7 +8799,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);
@@ -8815,7 +8815,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);
@@ -8827,7 +8827,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;
@@ -8848,14 +8848,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.
@@ -8887,7 +8886,7 @@ class MappableExprsHandler {
if (IsDevAddr) {
CI->ForDeviceAddr = true;
CI->ReturnDevicePointer = true;
- CI->FbNullify = FbNullify;
+ CI->HasUdpFbNullify = HasUdpFbNullify;
Found = true;
break;
} else {
@@ -8904,7 +8903,7 @@ class MappableExprsHandler {
VD == cast<DeclRefExpr>(AttachPtrExpr)->getDecl())) {
CI->ForDeviceAddr = IsDevAddr;
CI->ReturnDevicePointer = true;
- CI->FbNullify = FbNullify;
+ CI->HasUdpFbNullify = HasUdpFbNullify;
Found = true;
break;
}
@@ -8926,8 +8925,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);
@@ -8947,10 +8946,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);
}
}
@@ -9087,29 +9086,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);
}
}
More information about the llvm-branch-commits
mailing list