From a4743eba3c13799e667777764554fd44ccc6a33a Mon Sep 17 00:00:00 2001 From: Jennifer Yu Date: Mon, 4 Oct 2021 16:54:47 -0700 Subject: [PATCH] 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 --- clang/lib/CodeGen/CGOpenMPRuntime.cpp | 32 +++++++++++ .../test/OpenMP/nvptx_lambda_pointer_capturing.cpp | 64 ++++++++++++++++++++++ 2 files changed, 96 insertions(+) diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index e51d196..bf72897 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7481,6 +7481,9 @@ private: SmallVector> DevPointersMap; + /// Map between lambda declarations and their map type. + llvm::DenseMap 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()) 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()) { + 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. diff --git a/clang/test/OpenMP/nvptx_lambda_pointer_capturing.cpp b/clang/test/OpenMP/nvptx_lambda_pointer_capturing.cpp index 544158b..7c7e272 100644 --- a/clang/test/OpenMP/nvptx_lambda_pointer_capturing.cpp +++ b/clang/test/OpenMP/nvptx_lambda_pointer_capturing.cpp @@ -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 void omp_loop(int start, int end, F body) { @@ -24,6 +28,34 @@ void omp_loop(int start, int end, F body) { } } +template +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 +struct C { + static void xoo(const FTy& f) { + int x = 10; + #pragma omp target map(to:f) + f(x); + } +}; + +template +void zoo(const FTy &functor) { + C::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 -- 2.7.4