https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/210214
>From b705dab0717b6ef63b486341a6ef688029ea688b Mon Sep 17 00:00:00 2001 From: Abhinav Gaba <[email protected]> Date: Fri, 24 Jul 2026 11:36:17 -0700 Subject: [PATCH] [OpenMP] Propagate PRESENT to pointee entries in mapper codegen Extend map-type-modifier propagation in emitUserDefinedMapper to the PRESENT modifier, but only for entries that have an attach ptr (the pointee data, whose storage differs from the struct's own). A present modifier on the outer clause must require that pointee to be present on the device. This is gated on a new PropagatePresentToPointee argument, which Clang sets from CGM.getLangOpts().OpenMP >= 60. Before 6.0 the present modifier is treated as not applying to the pointee: the spec committee confirmed the divergence between the present motion modifier (to/from) and the present map-type modifier (map) was unintentional, to be fixed as an OpenMP 6.0 erratum. Only propagation is gated; present written directly in a mapper's own clause applies at all versions. A TODO notes PRESENT should also propagate to the struct's own members, which is blocked while pointer members use PTR_AND_OBJ. Update the present-check tests to their final 6.0-gated behavior. Co-Authored-By: Claude Opus 4.8 <[email protected]> --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 15 ++++-- ...t_map_nested_ptr_member_mapper_codegen.cpp | 15 +++--- .../llvm/Frontend/OpenMP/OMPIRBuilder.h | 17 +++++-- llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp | 47 ++++++++++++++----- .../mapper_map_mbr_then_present_mbr_ptee.c | 27 ++++++----- .../mapper_target_update_present_ptee.c | 27 ++++------- 6 files changed, 91 insertions(+), 57 deletions(-) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 87f409e2056d2..8e03d00e255fe 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -10595,8 +10595,11 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) { /// // Map-type-modifying bits (ALWAYS, DELETE, CLOSE) from the outer map /// // clause are propagated to each component, except ATTACH entries /// // (ATTACH|ALWAYS is reserved for attach(always), and other modifier -/// // bits have no meaning for ATTACH). PRESENT is handled separately. -/// imported_modifier_bits = type & (ALWAYS | DELETE | CLOSE); +/// // bits have no meaning for ATTACH). PRESENT is additionally +/// // propagated to pointee (attach-ptr) components at OpenMP >= 6.0. +/// present_bit = (v60 && c.hasAttachPtr()) ? PRESENT : 0; +/// imported_modifier_bits = +/// type & (ALWAYS | DELETE | CLOSE | present_bit); /// effective_type = c.isAttach() ? member_type /// : member_type | imported_modifier_bits; /// if (c.hasMapper()) @@ -10676,8 +10679,14 @@ void CGOpenMPRuntime::emitUserDefinedMapper(const OMPDeclareMapperDecl *D, CGM.getCXXABI().getMangleContext().mangleCanonicalTypeName(Ty, Out); std::string Name = getName({"omp_mapper", TyStr, D->getName()}); + // Propagate the PRESENT modifier to pointee (attach-ptr) entries only for + // OpenMP >= 6.0; before 6.0 the present modifier does not apply to the + // pointee (see the OpenMP 6.0 erratum on the present motion vs. map-type + // modifier divergence). + bool PropagatePresentToPointee = CGM.getLangOpts().OpenMP >= 60; llvm::Function *NewFn = cantFail(OMPBuilder.emitUserDefinedMapper( - PrivatizeAndGenMapInfoCB, ElemTy, Name, CustomMapperCB)); + PrivatizeAndGenMapInfoCB, ElemTy, Name, CustomMapperCB, + /*PreserveMemberOfFlags=*/false, PropagatePresentToPointee)); UDMMap.try_emplace(D, NewFn); if (CGF) FunctionUDMMap[CGF->CurFn].push_back(D); diff --git a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp index 821b5fd1652c4..40f22beb9bd7d 100644 --- a/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp +++ b/clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp @@ -3,12 +3,9 @@ // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -// PRESENT is propagated to pointee (attach-ptr) entries only at OpenMP >= 6.0. -// FIXME: that propagation is done in a follow-on; until then the CHECK-60 -// output below is identical to CHECK (the pointee entries carry map-type mask -// 1036 = ALWAYS|DELETE|CLOSE, without PRESENT). Once PRESENT is propagated, the -// attach-ptr pointee entries should use mask 5132 = ALWAYS|DELETE|CLOSE|PRESENT -// at 6.0. +// PRESENT is propagated to pointee (attach-ptr) entries only at OpenMP >= 6.0: +// at 6.0 those entries carry map-type mask 5132 = ALWAYS|DELETE|CLOSE|PRESENT, +// while the default (<= 5.2) CHECK uses 1036 = ALWAYS|DELETE|CLOSE (no PRESENT). // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK-60 // expected-no-diagnostics @@ -353,7 +350,7 @@ void foo(S2 *arr) { // CHECK-60: br label [[OMP_TYPE_END9]] // CHECK-60: omp.type.end9: // CHECK-60: [[OMP_MAPTYPE10:%.*]] = phi i64 [ 0, [[OMP_TYPE_ALLOC4]] ], [ 0, [[OMP_TYPE_TO6]] ], [ 0, [[OMP_TYPE_FROM8]] ], [ 0, [[OMP_TYPE_TO_ELSE7]] ] -// CHECK-60: [[TMP38:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-60: [[TMP38:%.*]] = and i64 [[TMP4]], 5132 // CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS11:%.*]] = or i64 [[OMP_MAPTYPE10]], [[TMP38]] // CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP15]], ptr [[X]], i64 [[TMP22]], i64 [[OMP_MAPTYPE_WITH_MODIFIERS11]], ptr null) // CHECK-60: [[TMP39:%.*]] = add nuw i64 562949953421315, [[TMP24]] @@ -377,7 +374,7 @@ void foo(S2 *arr) { // CHECK-60: br label [[OMP_TYPE_END17]] // CHECK-60: omp.type.end17: // CHECK-60: [[OMP_MAPTYPE18:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC12]] ], [ [[TMP44]], [[OMP_TYPE_TO14]] ], [ [[TMP46]], [[OMP_TYPE_FROM16]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE15]] ] -// CHECK-60: [[TMP47:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-60: [[TMP47:%.*]] = and i64 [[TMP4]], 5132 // CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS19:%.*]] = or i64 [[OMP_MAPTYPE18]], [[TMP47]] // CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP15]], ptr [[X]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS19]], ptr null) // CHECK-60: [[TMP48:%.*]] = add nuw i64 562949953421315, [[TMP24]] @@ -401,7 +398,7 @@ void foo(S2 *arr) { // CHECK-60: br label [[OMP_TYPE_END25]] // CHECK-60: omp.type.end25: // CHECK-60: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP51]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP53]], [[OMP_TYPE_TO22]] ], [ [[TMP55]], [[OMP_TYPE_FROM24]] ], [ [[TMP48]], [[OMP_TYPE_TO_ELSE23]] ] -// CHECK-60: [[TMP56:%.*]] = and i64 [[TMP4]], 1036 +// CHECK-60: [[TMP56:%.*]] = and i64 [[TMP4]], 5132 // CHECK-60: [[OMP_MAPTYPE_WITH_MODIFIERS27:%.*]] = or i64 [[OMP_MAPTYPE26]], [[TMP56]] // CHECK-60: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP17]], ptr [[Y]], i64 4, i64 [[OMP_MAPTYPE_WITH_MODIFIERS27]], ptr null) // CHECK-60: [[TMP57:%.*]] = and i64 [[TMP4]], 3 diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h index 33dc0e6c36774..2e99c084cfc97 100644 --- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h @@ -3624,8 +3624,13 @@ class OpenMPIRBuilder { /// // Map-type-modifying bits (ALWAYS, DELETE, CLOSE) from the outer /// // map clause are propagated to each component, except ATTACH /// // entries (ATTACH|ALWAYS is reserved for attach(always), and other - /// // modifier bits have no meaning for ATTACH). - /// imported_modifier_bits = type & (ALWAYS | DELETE | CLOSE); + /// // modifier bits have no meaning for ATTACH). PRESENT is + /// // additionally propagated to pointee (attach-ptr) components when + /// // PropagatePresentToPointee is set (OpenMP >= 6.0). + /// present_bit = (PropagatePresentToPointee && c.hasAttachPtr()) + /// ? PRESENT : 0; + /// imported_modifier_bits = type & (ALWAYS | DELETE | CLOSE | + /// present_bit); /// effective_type = c.isAttach() ? c.arg_type /// : c.arg_type | imported_modifier_bits; /// if (c.hasMapper()) @@ -3650,13 +3655,17 @@ class OpenMPIRBuilder { /// \param FuncName Optional param to specify mapper function name. /// \param CustomMapperCB Optional callback to generate code related to /// custom mappers. + /// \param PropagatePresentToPointee If true, the PRESENT map-type modifier + /// from the outer clause is propagated to pointee (attach-ptr) entries the + /// mapper inserts. Callers set this only for OpenMP >= 6.0; at earlier + /// versions the present modifier is treated as not applying to the pointee. LLVM_ABI Expected<Function *> emitUserDefinedMapper( function_ref<MapInfosOrErrorTy( InsertPointTy CodeGenIP, llvm::Value *PtrPHI, llvm::Value *BeginArg)> PrivAndGenMapInfoCB, llvm::Type *ElemTy, StringRef FuncName, - CustomMapperCallbackTy CustomMapperCB, - bool PreserveMemberOfFlags = false); + CustomMapperCallbackTy CustomMapperCB, bool PreserveMemberOfFlags = false, + bool PropagatePresentToPointee = false); /// Generator for '#omp target data' /// diff --git a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp index 63d03f3c446ba..36ccea21a784d 100644 --- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp @@ -10369,7 +10369,7 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper( llvm::Value *BeginArg)> GenMapInfoCB, Type *ElemTy, StringRef FuncName, CustomMapperCallbackTy CustomMapperCB, - bool PreserveMemberOfFlags) { + bool PreserveMemberOfFlags, bool PropagatePresentToPointee) { SmallVector<Type *> Params; Params.emplace_back(Builder.getPtrTy()); Params.emplace_back(Builder.getPtrTy()); @@ -10616,16 +10616,41 @@ Expected<Function *> OpenMPIRBuilder::emitUserDefinedMapper( // specified in the declared mapper. // // Map-type-modifying bits: ALWAYS, DELETE, CLOSE, PRESENT. - // TODO: PRESENT is not propagated here yet. Doing so requires - // distinguishing pointee entries from the struct's own storage; it is - // handled in a follow-on. - Value *ImportedModifierBits = Builder.CreateAnd( - MapType, - Builder.getInt64( - static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( - OpenMPOffloadMappingFlags::OMP_MAP_ALWAYS | - OpenMPOffloadMappingFlags::OMP_MAP_DELETE | - OpenMPOffloadMappingFlags::OMP_MAP_CLOSE))); + // + // ALWAYS/DELETE/CLOSE are propagated to every (non-ATTACH) entry. + // + // PRESENT is propagated only to entries that have an attach ptr + // (HasAttachPtr): the pointee data, which occupies a different storage + // block than the struct being mapped and so is not covered by the + // present-check on the struct's own storage. A present modifier on the + // outer clause must still require that pointee to be present on the device. + // + // This is gated on \p PropagatePresentToPointee (set by callers only for + // OpenMP >= 6.0). Before 6.0 the present modifier is treated as not + // applying to the pointee: the spec committee confirmed the divergence + // between the present "motion" modifier (to/from) and the present map-type + // modifier (map) was unintentional, to be fixed as an OpenMP 6.0 erratum, + // so for 5.2 present is ignored for the pointee for both map and to/from. + // + // TODO: PRESENT should also be propagated to the struct's own members + // (e.g. the s.x, s.y of map(present, mapper(id): s)) so that an absent + // member triggers the present-check. We cannot do that yet: while pointer + // members are mapped with PTR_AND_OBJ, a single combined entry allocates + // the whole struct (including the pointer's storage), so propagating + // PRESENT to it would wrongly require the pointer's pointee to be present. + // Enable member propagation once Clang stops emitting PTR_AND_OBJ and uses + // attach-style maps throughout. + uint64_t ModifierBits = + static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( + OpenMPOffloadMappingFlags::OMP_MAP_ALWAYS | + OpenMPOffloadMappingFlags::OMP_MAP_DELETE | + OpenMPOffloadMappingFlags::OMP_MAP_CLOSE); + if (PropagatePresentToPointee && Info->HasAttachPtr[I]) + ModifierBits |= + static_cast<std::underlying_type_t<OpenMPOffloadMappingFlags>>( + OpenMPOffloadMappingFlags::OMP_MAP_PRESENT); + Value *ImportedModifierBits = + Builder.CreateAnd(MapType, Builder.getInt64(ModifierBits)); Value *CurMapTypeWithModifiers = Builder.CreateOr( CurMapType, ImportedModifierBits, "omp.maptype.with.modifiers"); diff --git a/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c b/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c index bb8b6aa3c3e76..8d44e27146100 100644 --- a/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c +++ b/offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c @@ -1,18 +1,16 @@ // The mapper maps a struct member (s.x) and a pointee (s.p[0:10]). We pre-map // only s.x, then do map(present) on the mapper. The pointee s.p[0:10] is not -// present, so once PRESENT is propagated to the pointee (a follow-on, at OpenMP -// >= 6.0) the check must fail; at <= 5.2 present is not propagated, so it -// passes. -// -// FIXME: PRESENT is not propagated to the pointee yet, so the run currently -// completes ("done") at BOTH versions. Once it is propagated: -// EXPECTED (5.2): the run completes ("done"). -// EXPECTED (6.0): the present check fails for the absent pointee s1.p[0:10]. +// present, so the propagated present modifier must fail the check -- but only +// at OpenMP >= 6.0, since present is not propagated to the pointee before then. + +// OpenMP <= 5.2: present is not propagated to the pointee, so the run succeeds. // RUN: %libomptarget-compile-generic -fopenmp-version=52 // RUN: %libomptarget-run-generic 2>&1 \ // RUN: | %fcheck-generic --check-prefixes=CHECK,CHECK-52 + +// OpenMP 6.0: present is propagated to the pointee; the check fails. // RUN: %libomptarget-compile-generic -fopenmp-version=60 -// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: %libomptarget-run-fail-generic 2>&1 \ // RUN: | %fcheck-generic --check-prefixes=CHECK,CHECK-60 #include <omp.h> @@ -48,11 +46,14 @@ int main() { print_status(&s1.p[0], "p[0]"); // CHECK: p[0] is not present #pragma omp target enter data map(present, alloc : s1) - // Once PRESENT is propagated to the pointee, at 5.2 the run completes past - // this point; at 6.0 the present check on the absent pointee s1.p[0:10] - // fails here. + // At 5.2 the run completes past this point; at 6.0 the present check on the + // absent pointee s1.p[0:10] fails here. // CHECK-52: done - // CHECK-60: done + // + // clang-format off + // CHECK-60: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes) + // CHECK-60: omptarget fatal error 1: failure of target construct while offloading is mandatory + // clang-format on fprintf(stderr, "done\n"); } diff --git a/offload/test/mapping/mapper_target_update_present_ptee.c b/offload/test/mapping/mapper_target_update_present_ptee.c index 10f11392ed76c..f241be882b6b7 100644 --- a/offload/test/mapping/mapper_target_update_present_ptee.c +++ b/offload/test/mapping/mapper_target_update_present_ptee.c @@ -6,17 +6,14 @@ // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=CHECK-60 // Out-of-bounds: s.p[0:20] extends beyond the mapped region x[0:10]. At OpenMP -// 6.0 the present modifier should be propagated to the pointee s.p[0:20], so -// this should run-fail with a present error; at <= 5.2 present is (correctly) -// not applied to the pointee, so the run succeeds. -// FIXME: the present modifier is not yet propagated to the pointee, so the OOB -// run currently succeeds at 6.0 too; propagating it is done in a follow-on. -// EXPECTED (6.0): run-fail with a present-modifier error for s.p[0:20]. +// 6.0 the present modifier is propagated to the pointee s.p[0:20], so the +// update fails the present check; at <= 5.2 present is not applied to the +// pointee, so it succeeds. // RUN: %libomptarget-compile-generic -fopenmp-version=52 -DOUT_OF_BOUNDS // RUN: %libomptarget-run-generic 2>&1 \ // RUN: | %fcheck-generic --check-prefix=CHECK-52-OOB // RUN: %libomptarget-compile-generic -fopenmp-version=60 -DOUT_OF_BOUNDS -// RUN: %libomptarget-run-generic 2>&1 \ +// RUN: %libomptarget-run-fail-generic 2>&1 \ // RUN: | %fcheck-generic --check-prefix=CHECK-60-OOB #include <stdio.h> @@ -54,6 +51,11 @@ int main() { // CHECK-60-OOB: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]] fprintf(stderr, "addr=%p, size=%zu\n", &s.p[0], 20 * sizeof(s.p[0])); + // At 6.0 the out-of-bounds pointee fails the present check inside f1(). + // clang-format off + // CHECK-60-OOB: omptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes) + // CHECK-60-OOB: omptarget fatal error 1: failure of target construct while offloading is mandatory + // clang-format on #pragma omp target data map(from : s.y, x) { f1(); @@ -64,21 +66,12 @@ int main() { // CHECK-60: 333 333 fprintf(stderr, "%d %d\n", x[0], s.y); - // Out-of-bounds: present is not yet propagated to the pointee, so both - // versions currently complete past the update. - // - // 5.2: present is (correctly) never applied to the pointee, so the update + // 5.2 out-of-bounds: present is not applied to the pointee, so the update // completes. // FIXME: even at 5.2, the update should still have happened for the subset of // the pointee that is present (x[0:10]), instead of being silently ignored // (x[0] currently reads back as garbage). Once that is fixed: // EXPECTED-52-OOB: 333 333 // CHECK-52-OOB: done - // - // 6.0: present should be propagated to the pointee, so the update should fail - // the present check for s.p[0:20]. That propagation is done in a follow-on; - // until then the run completes. - // EXPECTED-60-OOB: run-fail with a present-modifier error for s.p[0:20]. - // CHECK-60-OOB: done fprintf(stderr, "done\n"); } _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
