[llvm-branch-commits] [clang] [llvm] [OpenMP] Propagate PRESENT to pointee entries in mapper codegen (PR #210214)
Abhinav Gaba via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Aug 3 16:49:26 PDT 2026
https://github.com/abhinavgaba updated https://github.com/llvm/llvm-project/pull/210214
>From 9dfd09764d54e970c772fdbe17e19f112fb78301 Mon Sep 17 00:00:00 2001
From: Abhinav Gaba <abhinav.gaba at intel.com>
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 <noreply at anthropic.com>
---
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 6edabd4e42d88..3a4b5a351801a 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 2fd98feea0c33..3e46f74e71d74 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -3636,8 +3636,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())
@@ -3662,13 +3667,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 3df035f75746c..cb3cdcb601bc3 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10426,7 +10426,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());
@@ -10673,16 +10673,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");
}
More information about the llvm-branch-commits
mailing list