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