llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-offload
Author: Abhinav Gaba (abhinavgaba)
<details>
<summary>Changes</summary>
For cases like:
```c
#pragma omp declare_mapper (default: S s) map(s.x, s.p[0:10])
S s1;
...
#pragma omp target_enter_data map(present, alloc: s1)
```
After "mapper-expansion", the behavior of the above should be equivalent to:
```
#pragma omp target_enter_data map(present, alloc: s1.x) map(present, alloc:
s.p[0:10])
```
i.e. The `present` map-type needs to be propagated to the map for the pointee
of `s.p`. That was not happening prior to this change.
We try to limit this to entries that have their own "attach-ptr" (like
`s.p[0:10]`) and thus occupy a separate storage block than the base variable
for which the mapper is declared (like `s`).
Note that we do this only with OpenMP 6.0+, since the wording in 5.2 implies
that PRESENT applies only to `s.x` after mapper-expansion.
TODO: `PRESENT` bit should also be propagated to non-pointee entries resulting
from mapper-expansion, e.g.
```c
struct S {
int x; int y; int *p;
};
S s1;
#pragma omp declare_mapper map(s.x, s.y)
...
#pragma omp target_enter_data map(present: s1)
// present has to be propagated inwards to `s.x, s.y` so that it doesn't
// assert that the full `s1` is "present".
```
However we cannot do that until we have removed the remaining uses of
PTR_AND_OBJ maps map-type (which is used for byrefs in clang at the moment).
---
Full diff: https://github.com/llvm/llvm-project/pull/210214.diff
6 Files Affected:
- (modified) clang/lib/CodeGen/CGOpenMPRuntime.cpp (+13-3)
- (modified) clang/test/OpenMP/target_map_nested_ptr_member_mapper_codegen.cpp
(+6-9)
- (modified) llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h (+15-4)
- (modified) llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp (+36-11)
- (modified) offload/test/mapping/mapper_map_mbr_then_present_mbr_ptee.c
(+14-13)
- (modified) offload/test/mapping/mapper_target_update_present_ptee.c (+10-17)
``````````diff
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 6edabd4e42d88..51f1e8b22fcfc 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -10595,8 +10595,12 @@ 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 components with HasAttachPtr (the pointee data) 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 +10680,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 the pointee entries (those with
+ // HasAttachPtr) 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..18d5133098162 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 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 fc0aaeafc4fcf..1965f7b983805 100644
--- a/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
+++ b/llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h
@@ -3693,8 +3693,14 @@ 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 components with HasAttachPtr (the
+ /// // pointee data) 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())
@@ -3719,13 +3725,18 @@ 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 the pointee entries the mapper
+ /// inserts, i.e. those with HasAttachPtr. 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 db55fc3ab0fa0..c85cfe15d058c 100644
--- a/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+++ b/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
@@ -10542,7 +10542,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());
@@ -10801,16 +10801,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..3e1856c2567ce 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: message: device mapping required by 'present' map type modifier
does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
+ // CHECK-60: 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..a3adcd6eb196a 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: message: device mapping required by 'present' motion
modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]]
bytes)
+ // CHECK-60-OOB: 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");
}
``````````
</details>
https://github.com/llvm/llvm-project/pull/210214
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
