forked from OSchip/llvm-project
Fix assert of "Unable to find base lambda address" from
adjustMemberOfForLambdaCaptures. The problem is happening when user passes lambda function with reference type in the map clause. The natural of the problem when processing generateInfoForCapture, the BasePointer is generated with new load for a lambda variable with reference type. It is not expected in adjustMemberOfForLambdaCaptures. One way to fix this is to skipping call to generateInfoForCapture for map(to:lambda). The map info will be generated later in the call to generateDefaultMapInfo samiler as firsprivate clase. This to fix https://bugs.llvm.org/show_bug.cgi?id=52071 Differential Revision:https://reviews.llvm.org/D111115
This commit is contained in:
parent
c7e72784ae
commit
a4743eba3c
|
@ -7481,6 +7481,9 @@ private:
|
|||
SmallVector<OMPClauseMappableExprCommon::MappableExprComponentListRef, 4>>
|
||||
DevPointersMap;
|
||||
|
||||
/// Map between lambda declarations and their map type.
|
||||
llvm::DenseMap<const ValueDecl *, const OMPMapClause *> LambdasMap;
|
||||
|
||||
llvm::Value *getExprTypeSize(const Expr *E) const {
|
||||
QualType ExprTy = E->getType().getCanonicalType();
|
||||
|
||||
|
@ -8442,6 +8445,15 @@ private:
|
|||
return MappableExprsHandler::OMP_MAP_PRIVATE |
|
||||
MappableExprsHandler::OMP_MAP_TO;
|
||||
}
|
||||
auto I = LambdasMap.find(Cap.getCapturedVar()->getCanonicalDecl());
|
||||
if (I != LambdasMap.end())
|
||||
// for map(to: lambda): using user specified map type.
|
||||
return getMapTypeBits(
|
||||
I->getSecond()->getMapType(), I->getSecond()->getMapTypeModifiers(),
|
||||
/*MotionModifiers=*/llvm::None, I->getSecond()->isImplicit(),
|
||||
/*AddPtrFlag=*/false,
|
||||
/*AddIsTargetParamFlag=*/false,
|
||||
/*isNonContiguous=*/false);
|
||||
return MappableExprsHandler::OMP_MAP_TO |
|
||||
MappableExprsHandler::OMP_MAP_FROM;
|
||||
}
|
||||
|
@ -8906,6 +8918,21 @@ public:
|
|||
for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
|
||||
for (auto L : C->component_lists())
|
||||
DevPointersMap[std::get<0>(L)].push_back(std::get<1>(L));
|
||||
// Extract map information.
|
||||
for (const auto *C : Dir.getClausesOfKind<OMPMapClause>()) {
|
||||
if (C->getMapType() != OMPC_MAP_to)
|
||||
continue;
|
||||
for (auto L : C->component_lists()) {
|
||||
const ValueDecl *VD = std::get<0>(L);
|
||||
const auto *RD = VD ? VD->getType()
|
||||
.getCanonicalType()
|
||||
.getNonReferenceType()
|
||||
->getAsCXXRecordDecl()
|
||||
: nullptr;
|
||||
if (RD && RD->isLambda())
|
||||
LambdasMap.try_emplace(std::get<0>(L), C);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/// Constructor for the declare mapper directive.
|
||||
|
@ -9118,6 +9145,11 @@ public:
|
|||
? nullptr
|
||||
: Cap->getCapturedVar()->getCanonicalDecl();
|
||||
|
||||
// for map(to: lambda): skip here, processing it in
|
||||
// generateDefaultMapInfo
|
||||
if (LambdasMap.count(VD))
|
||||
return;
|
||||
|
||||
// If this declaration appears in a is_device_ptr clause we just have to
|
||||
// pass the pointer by value. If it is a reference to a declaration, we just
|
||||
// pass its value.
|
||||
|
|
|
@ -15,6 +15,10 @@
|
|||
// CHECK-DAG: [[TYPES_TEMPLATE:@.+]] = private {{.+}} constant [5 x i64] [i64 800, i64 800, i64 673, i64 844424930132752, i64 844424930132752]
|
||||
// CHECK-DAG: [[SIZES:@.+]] = private {{.+}} constant [3 x i[[PTRSZ:32|64]]] [i{{32|64}} {{8|16}}, i{{32|64}} 0, i{{32|64}} 0]
|
||||
// CHECK-DAG: [[TYPES:@.+]] = private {{.+}} constant [3 x i64] [i64 673, i64 281474976711440, i64 281474976711440]
|
||||
// CHECK-DAG: [[TYPES3:@.+]] = private {{.+}} constant [3 x i64] [i64 545, i64 281474976711440, i64 800]
|
||||
// CHECK-DAG: [[TYPES11:@.+]] = private {{.+}} constant [5 x i64] [i64 800, i64 800, i64 549, i64 844424930132752, i64 844424930132752]
|
||||
// CHECK-DAG: [[TYPES13:@.+]] = private {{.+}} constant [2 x i64] [i64 545, i64 281474976711440]
|
||||
// CHECK-DAG: [[TYPES15:@.+]] = private {{.+}} constant [2 x i64] [i64 673, i64 281474976711440]
|
||||
|
||||
template <typename F>
|
||||
void omp_loop(int start, int end, F body) {
|
||||
|
@ -24,6 +28,34 @@ void omp_loop(int start, int end, F body) {
|
|||
}
|
||||
}
|
||||
|
||||
template <typename F>
|
||||
void omp_loop_ref(int start, int end, F& body) {
|
||||
#pragma omp target teams distribute parallel for map(always, to: body)
|
||||
for (int i = start; i < end; ++i) {
|
||||
body(i);
|
||||
}
|
||||
int *p;
|
||||
const auto &body_ref = [=](int i) {p[i]=0;};
|
||||
#pragma omp target map(to: body_ref)
|
||||
body_ref(10);
|
||||
#pragma omp target
|
||||
body_ref(10);
|
||||
}
|
||||
|
||||
template <class FTy>
|
||||
struct C {
|
||||
static void xoo(const FTy& f) {
|
||||
int x = 10;
|
||||
#pragma omp target map(to:f)
|
||||
f(x);
|
||||
}
|
||||
};
|
||||
|
||||
template <class FTy>
|
||||
void zoo(const FTy &functor) {
|
||||
C<FTy>::xoo(functor);
|
||||
}
|
||||
|
||||
// CHECK: define {{.*}}[[MAIN:@.+]](
|
||||
int main()
|
||||
{
|
||||
|
@ -32,6 +64,7 @@ int main()
|
|||
auto body = [=](int i){
|
||||
p[i] = q[i];
|
||||
};
|
||||
zoo([=](int i){p[i] = 0;});
|
||||
|
||||
#pragma omp target teams distribute parallel for
|
||||
for (int i = 0; i < 100; ++i) {
|
||||
|
@ -82,6 +115,7 @@ int main()
|
|||
|
||||
|
||||
omp_loop(0,100,body);
|
||||
omp_loop_ref(0,100,body);
|
||||
}
|
||||
|
||||
// CHECK: [[BASE_PTRS:%.+]] = alloca [5 x i8*]{{.+}}
|
||||
|
@ -122,4 +156,34 @@ int main()
|
|||
// CHECK: [[PTRS_GEP:%.+]] = getelementptr {{.+}} [5 x {{.+}}*], [5 x {{.+}}*]* [[PTRS]], {{.+}} 0, {{.+}} 0
|
||||
// CHECK: {{%.+}} = call{{.+}} @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}, {{.+}}, {{.+}}, i8** [[BASES_GEP]], i8** [[PTRS_GEP]], i[[PTRSZ]]* getelementptr inbounds ([5 x i{{.+}}], [5 x i{{.+}}]* [[SIZES_TEMPLATE]], i{{.+}} 0, i{{.+}} 0), i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[TYPES_TEMPLATE]], i{{.+}} 0, i{{.+}} 0), i8** null, i8** null, {{.+}}, {{.+}})
|
||||
|
||||
// CHECK: define internal void @{{.+}}omp_loop_ref{{.+}}(
|
||||
// CHECK: [[BODY:%body.addr]] = alloca %class.anon*
|
||||
// CHECK: [[TMP:%tmp]] = alloca %class.anon*
|
||||
// CHECK: [[BODY_REF:%body_ref]] = alloca %class.anon.1*
|
||||
// CHECK: [[REF_TMP:%ref.tmp]] = alloca %class.anon.1
|
||||
// CHECK: [[TMP8:%tmp.+]] = alloca %class.anon.1*
|
||||
// CHECK: [[L0:%.+]] = load %class.anon*, %class.anon** [[BODY]]
|
||||
// CHECK: store %class.anon* [[L0]], %class.anon** [[TMP]]
|
||||
// CHECK: [[L5:%.+]] = load %class.anon*, %class.anon** [[TMP]]
|
||||
// CHECK-NOT [[L6:%.+]] = load %class.anon*, %class.anon** [[TMP]]
|
||||
// CHECK-NOT [[L7:%.+]] = load %class.anon*, %class.anon** [[TMP]]
|
||||
// CHECK: store %class.anon.1* [[REF_TMP]], %class.anon.1** [[BODY_REF]]
|
||||
// CHECK:[[L47:%.+]] = load %class.anon.1*, %class.anon.1** [[BODY_REF]]
|
||||
// CHECK: store %class.anon.1* [[L47]], %class.anon.1** [[TMP8]]
|
||||
// CHECK: [[L48:%.+]] = load %class.anon.1*, %class.anon.1** [[TMP8]]
|
||||
// CHECK-NOT: [[L49:%.+]] = load %class.anon.1*, %class.anon.1** [[TMP8]]
|
||||
// CHECK-NOT: [[L50:%.+]] = load %class.anon.1*, %class.anon.1** [[TMP8]]
|
||||
// CHECK: ret void
|
||||
|
||||
// CHECK: define internal void @{{.+}}xoo{{.+}}(
|
||||
// CHECK: [[FADDR:%f.addr]] = alloca %class.anon.0*
|
||||
// CHECK: [[L0:%.+]] = load %class.anon.0*, %class.anon.0** [[FADDR]]
|
||||
// CHECK: store %class.anon.0* [[L0]], %class.anon.0** [[TMP:%tmp]]
|
||||
// CHECK: [[L1:%.+]] = load %class.anon.0*, %class.anon.0** [[TMP]]
|
||||
// CHECK-NOT: %4 = load %class.anon.0*, %class.anon.0** [[TMP]]
|
||||
// CHECK-NOT: %5 = load %class.anon.0*, %class.anon.0** [[TMP]]
|
||||
// CHECK: [[L4:%.+]] = getelementptr inbounds %class.anon.0, %class.anon.0* [[L1]], i32 0, i32 0
|
||||
// CHECK: [[L5:%.+]] = load i{{.*}}*, i{{.*}}** [[L4]]
|
||||
// CHECK: ret void
|
||||
|
||||
#endif
|
||||
|
|
Loading…
Reference in New Issue