[clang] a4743eb - Fix assert of "Unable to find base lambda address" from

Jennifer Yu via cfe-commits cfe-commits at lists.llvm.org
Wed Oct 6 14:38:49 PDT 2021


Author: Jennifer Yu
Date: 2021-10-06T14:14:28-07:00
New Revision: a4743eba3c13799e667777764554fd44ccc6a33a

URL: https://github.com/llvm/llvm-project/commit/a4743eba3c13799e667777764554fd44ccc6a33a
DIFF: https://github.com/llvm/llvm-project/commit/a4743eba3c13799e667777764554fd44ccc6a33a.diff

LOG: 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

Added: 
    

Modified: 
    clang/lib/CodeGen/CGOpenMPRuntime.cpp
    clang/test/OpenMP/nvptx_lambda_pointer_capturing.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index e51d196b870f..bf72897d81b0 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7481,6 +7481,9 @@ class MappableExprsHandler {
       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 @@ class MappableExprsHandler {
       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 @@ class MappableExprsHandler {
     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 @@ class MappableExprsHandler {
                               ? 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 544158ba4594..7c7e27239680 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 <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


        


More information about the cfe-commits mailing list