jyu2 updated this revision to Diff 377421.
jyu2 added a comment.
Adding check for map type in test. Thanks Alexey.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D111115/new/
https://reviews.llvm.org/D111115
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/nvptx_lambda_pointer_capturing.cpp
Index: clang/test/OpenMP/nvptx_lambda_pointer_capturing.cpp
===================================================================
--- clang/test/OpenMP/nvptx_lambda_pointer_capturing.cpp
+++ 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 @@
}
}
+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 @@
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 @@
omp_loop(0,100,body);
+ omp_loop_ref(0,100,body);
}
// CHECK: [[BASE_PTRS:%.+]] = alloca [5 x i8*]{{.+}}
@@ -122,4 +156,34 @@
// 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
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7481,6 +7481,9 @@
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 @@
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 @@
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 @@
? 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.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
