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 <[email protected]> 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 <[email protected]> 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); } } _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
