From 4c2015e24d03913fedb0ac31ee6426dbfafb5e38 Mon Sep 17 00:00:00 2001 From: Alexey Bataev <a.bataev@hotmail.com> Date: Fri, 6 Oct 2017 16:17:25 +0000 Subject: [PATCH] [OPENMP] Capture references to global variables. In C++11 variable to global variables are considered as constant expressions and these variables are not captured in the outlined regions. Patch allows capturing of such variables in the OpenMP regions. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@315074 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/CodeGen/CGExpr.cpp | 8 +++-- lib/Sema/SemaExpr.cpp | 3 +- test/OpenMP/for_firstprivate_codegen.cpp | 30 ++++++++++++------- test/OpenMP/for_lastprivate_codegen.cpp | 13 ++++++-- test/OpenMP/for_private_codegen.cpp | 18 ++++++++--- .../teams_distribute_firstprivate_codegen.cpp | 28 +++++++++-------- .../teams_distribute_private_codegen.cpp | 7 +++-- 7 files changed, 73 insertions(+), 34 deletions(-) diff --git a/lib/CodeGen/CGExpr.cpp b/lib/CodeGen/CGExpr.cpp index 1bb57b6e2b..b1b7c1b60b 100644 --- a/lib/CodeGen/CGExpr.cpp +++ b/lib/CodeGen/CGExpr.cpp @@ -2297,8 +2297,12 @@ LValue CodeGenFunction::EmitDeclRefLValue(const DeclRefExpr *E) { VD->isUsableInConstantExpressions(getContext()) && VD->checkInitIsICE() && // Do not emit if it is private OpenMP variable. - !(E->refersToEnclosingVariableOrCapture() && CapturedStmtInfo && - LocalDeclMap.count(VD))) { + !(E->refersToEnclosingVariableOrCapture() && + ((CapturedStmtInfo && + (LocalDeclMap.count(VD->getCanonicalDecl()) || + CapturedStmtInfo->lookup(VD->getCanonicalDecl()))) || + LambdaCaptureFields.lookup(VD->getCanonicalDecl()) || + isa<BlockDecl>(CurCodeDecl)))) { llvm::Constant *Val = ConstantEmitter(*this).emitAbstract(E->getLocation(), *VD->evaluateValue(), diff --git a/lib/Sema/SemaExpr.cpp b/lib/Sema/SemaExpr.cpp index dbefcdb66e..4e27170ae9 100644 --- a/lib/Sema/SemaExpr.cpp +++ b/lib/Sema/SemaExpr.cpp @@ -14893,7 +14893,8 @@ static void DoMarkVarDeclReferenced(Sema &SemaRef, SourceLocation Loc, IsVariableAConstantExpression(Var, SemaRef.Context)) { // A reference initialized by a constant expression can never be // odr-used, so simply ignore it. - if (!Var->getType()->isReferenceType()) + if (!Var->getType()->isReferenceType() || + (SemaRef.LangOpts.OpenMP && SemaRef.IsOpenMPCapturedDecl(Var))) SemaRef.MaybeODRUseExprs.insert(E); } else if (OdrUseContext) { MarkVarDeclODRUsed(Var, Loc, SemaRef, diff --git a/test/OpenMP/for_firstprivate_codegen.cpp b/test/OpenMP/for_firstprivate_codegen.cpp index 1e1d404b48..590739203f 100644 --- a/test/OpenMP/for_firstprivate_codegen.cpp +++ b/test/OpenMP/for_firstprivate_codegen.cpp @@ -83,6 +83,7 @@ int main() { // LAMBDA: alloca i{{[0-9]+}}, // LAMBDA: [[G_PRIVATE_ADDR:%.+]] = alloca i{{[0-9]+}}, // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[G1_PRIVATE_REF:%.+]] = alloca i{{[0-9]+}}*, // LAMBDA: [[SIVAR2_PRIVATE_ADDR:%.+]] = alloca i{{[0-9]+}}, // LAMBDA: store i{{[0-9]+}}* [[SIVAR_REF]], i{{[0-9]+}}** %{{.+}}, @@ -91,20 +92,26 @@ int main() { // LAMBDA: [[G_VAL:%.+]] = load volatile i{{[0-9]+}}, i{{[0-9]+}}* [[G]] // LAMBDA: store i{{[0-9]+}} [[G_VAL]], i{{[0-9]+}}* [[G_PRIVATE_ADDR]] + // LAMBDA: store i{{[0-9]+}}* [[G1_PRIVATE_ADDR]], i{{[0-9]+}}** [[G1_PRIVATE_REF]], // LAMBDA: [[SIVAR2_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[SIVAR2_PRIVATE_ADDR_REF]] // LAMBDA: store i{{[0-9]+}} [[SIVAR2_VAL]], i{{[0-9]+}}* [[SIVAR2_PRIVATE_ADDR]] // LAMBDA-NOT: call void @__kmpc_barrier( g = 1; - g1 = 1; - sivar = 2; + g1 = 2; + sivar = 3; // LAMBDA: call void @__kmpc_for_static_init_4( // LAMBDA: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[G_PRIVATE_ADDR]], - // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[SIVAR2_PRIVATE_ADDR]], + // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G1_PRIVATE_REF]], + // LAMBDA: store volatile i{{[0-9]+}} 2, i{{[0-9]+}}* [[G1_PRIVATE_ADDR]], + // LAMBDA: store i{{[0-9]+}} 3, i{{[0-9]+}}* [[SIVAR2_PRIVATE_ADDR]], // LAMBDA: [[G_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // LAMBDA: store i{{[0-9]+}}* [[G_PRIVATE_ADDR]], i{{[0-9]+}}** [[G_PRIVATE_ADDR_REF]] - // LAMBDA: [[SIVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G1_PRIVATE_REF]], + // LAMBDA: store i{{[0-9]+}}* [[G1_PRIVATE_ADDR]], i{{[0-9]+}}** [[G1_PRIVATE_ADDR_REF]] + // LAMBDA: [[SIVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // LAMBDA: store i{{[0-9]+}}* [[SIVAR2_PRIVATE_ADDR]], i{{[0-9]+}}** [[SIVAR_PRIVATE_ADDR_REF]] // LAMBDA: call void [[INNER_LAMBDA:@.+]](%{{.+}}* [[ARG]]) // LAMBDA: call void @__kmpc_for_static_fini( @@ -112,17 +119,20 @@ int main() { [&]() { // LAMBDA: define {{.+}} void [[INNER_LAMBDA]](%{{.+}}* [[ARG_PTR:%.+]]) // LAMBDA: store %{{.+}}* [[ARG_PTR]], %{{.+}}** [[ARG_PTR_REF:%.+]], - g = 2; - g1 = 2; - sivar = 4; + g = 4; + g1 = 5; + sivar = 6; // LAMBDA: [[ARG_PTR:%.+]] = load %{{.+}}*, %{{.+}}** [[ARG_PTR_REF]] // LAMBDA: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // LAMBDA: [[G_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G_PTR_REF]] - // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[G_REF]] - // LAMBDA: [[SIVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[G_REF]] + // LAMBDA: [[G1_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G1_PTR_REF]] + // LAMBDA: store i{{[0-9]+}} 5, i{{[0-9]+}}* [[G1_REF]] + // LAMBDA: [[SIVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // LAMBDA: [[SIVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SIVAR_PTR_REF]] - // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SIVAR_REF]] + // LAMBDA: store i{{[0-9]+}} 6, i{{[0-9]+}}* [[SIVAR_REF]] }(); } }(); diff --git a/test/OpenMP/for_lastprivate_codegen.cpp b/test/OpenMP/for_lastprivate_codegen.cpp index f0fba042c5..47c24ed102 100644 --- a/test/OpenMP/for_lastprivate_codegen.cpp +++ b/test/OpenMP/for_lastprivate_codegen.cpp @@ -246,6 +246,7 @@ int main() { // LAMBDA: alloca i{{[0-9]+}}, // LAMBDA: [[G_PRIVATE_ADDR:%.+]] = alloca i{{[0-9]+}}, align 128 // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[G1_PRIVATE_REF:%.+]] = alloca i{{[0-9]+}}*, // LAMBDA: [[SIVAR_PRIVATE_ADDR:%.+]] = alloca i{{[0-9]+}}, // LAMBDA: [[SIVAR_PRIVATE_ADDR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** %{{.+}}, @@ -254,10 +255,15 @@ int main() { // LAMBDA: call {{.+}} @__kmpc_for_static_init_4(%{{.+}}* @{{.+}}, i32 [[GTID]], i32 34, i32* [[IS_LAST_ADDR:%.+]], i32* %{{.+}}, i32* %{{.+}}, i32* %{{.+}}, i32 1, i32 1) // LAMBDA: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[G_PRIVATE_ADDR]], + // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G1_PRIVATE_REF]], + // LAMBDA: store volatile i{{[0-9]+}} 1, i{{[0-9]+}}* [[G1_PRIVATE_ADDR]], // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[SIVAR_PRIVATE_ADDR]], // LAMBDA: [[G_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // LAMBDA: store i{{[0-9]+}}* [[G_PRIVATE_ADDR]], i{{[0-9]+}}** [[G_PRIVATE_ADDR_REF]] - // LAMBDA: [[SIVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G1_PRIVATE_REF]], + // LAMBDA: store i{{[0-9]+}}* [[G1_PRIVATE_ADDR]], i{{[0-9]+}}** [[G1_PRIVATE_ADDR_REF]] + // LAMBDA: [[SIVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // LAMBDA: store i{{[0-9]+}}* [[SIVAR_PRIVATE_ADDR]], i{{[0-9]+}}** [[SIVAR_PRIVATE_ADDR_REF]] // LAMBDA: call void [[INNER_LAMBDA:@.+]](%{{.+}}* [[ARG]]) // LAMBDA: call void @__kmpc_for_static_fini(%{{.+}}* @{{.+}}, i32 [[GTID]]) @@ -291,7 +297,10 @@ int main() { // LAMBDA: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // LAMBDA: [[G_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G_PTR_REF]] // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[G_REF]] - // LAMBDA: [[SIVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G1_PTR_REF]] + // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[G1_REF]] + // LAMBDA: [[SIVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // LAMBDA: [[SIVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SIVAR_PTR_REF]] // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SIVAR_REF]] }(); diff --git a/test/OpenMP/for_private_codegen.cpp b/test/OpenMP/for_private_codegen.cpp index 5bad8d0068..f1416d75b6 100644 --- a/test/OpenMP/for_private_codegen.cpp +++ b/test/OpenMP/for_private_codegen.cpp @@ -52,6 +52,8 @@ int main() { for (int i = 0; i < 2; ++i) { // LAMBDA: define{{.*}} internal{{.*}} void [[OMP_REGION]](i32* noalias %{{.+}}, i32* noalias %{{.+}}) // LAMBDA: [[G_PRIVATE_ADDR:%.+]] = alloca double, + // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = alloca double, + // LAMBDA: [[G1_PRIVATE_REF:%.+]] = alloca double*, // LAMBDA: [[SVAR_PRIVATE_ADDR:%.+]] = alloca i{{[0-9]+}}, // LAMBDA: [[SFVAR_PRIVATE_ADDR:%.+]] = alloca float, g = 1; @@ -60,13 +62,18 @@ int main() { sfvar = 4.0; // LAMBDA: call {{.*}}void @__kmpc_for_static_init_4( // LAMBDA: store double 1.0{{.+}}, double* [[G_PRIVATE_ADDR]], + // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = load double*, double** [[G1_PRIVATE_REF]], + // LAMBDA: store volatile double 1.0{{.+}}, double* [[G1_PRIVATE_ADDR]], // LAMBDA: store i{{[0-9]+}} 3, i{{[0-9]+}}* [[SVAR_PRIVATE_ADDR]], // LAMBDA: store float 4.0{{.+}}, float* [[SFVAR_PRIVATE_ADDR]], // LAMBDA: [[G_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // LAMBDA: store double* [[G_PRIVATE_ADDR]], double** [[G_PRIVATE_ADDR_REF]] - // LAMBDA: [[SVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PRIVATE_ADDR:%.+]] = load double*, double** [[G1_PRIVATE_REF]], + // LAMBDA: store double* [[G1_PRIVATE_ADDR]], double** [[G1_PRIVATE_ADDR_REF]] + // LAMBDA: [[SVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // LAMBDA: store i{{[0-9]+}}* [[SVAR_PRIVATE_ADDR]], i{{[0-9]+}}** [[SVAR_PRIVATE_ADDR_REF]] - // LAMBDA: [[SFVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // LAMBDA: [[SFVAR_PRIVATE_ADDR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG:%.+]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 // LAMBDA: store float* [[SFVAR_PRIVATE_ADDR]], float** [[SFVAR_PRIVATE_ADDR_REF]] // LAMBDA: call{{.*}} void [[INNER_LAMBDA:@.+]](%{{.+}}* [[ARG]]) // LAMBDA: call {{.*}}void @__kmpc_for_static_fini( @@ -81,10 +88,13 @@ int main() { // LAMBDA: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // LAMBDA: [[G_REF:%.+]] = load double*, double** [[G_PTR_REF]] // LAMBDA: store double 2.0{{.+}}, double* [[G_REF]] - // LAMBDA: [[SVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_REF:%.+]] = load double*, double** [[G1_PTR_REF]] + // LAMBDA: store double 2.0{{.+}}, double* [[G1_REF]] + // LAMBDA: [[SVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // LAMBDA: [[SVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SVAR_PTR_REF]] // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SVAR_REF]] - // LAMBDA: [[SFVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // LAMBDA: [[SFVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 // LAMBDA: [[SFVAR_REF:%.+]] = load float*, float** [[SFVAR_PTR_REF]] // LAMBDA: store float 8.0{{.+}}, float* [[SFVAR_REF]] }(); diff --git a/test/OpenMP/teams_distribute_firstprivate_codegen.cpp b/test/OpenMP/teams_distribute_firstprivate_codegen.cpp index 1f52ebea1a..aef288d11b 100644 --- a/test/OpenMP/teams_distribute_firstprivate_codegen.cpp +++ b/test/OpenMP/teams_distribute_firstprivate_codegen.cpp @@ -73,7 +73,7 @@ int main() { // LAMBDA: call void [[OUTER_LAMBDA:@.+]]( [&]() { // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // LAMBDA: call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0)) // LAMBDA: call void @[[LOFFL1:.+]](i{{64|32}} %{{.+}}) // LAMBDA: ret #pragma omp target @@ -84,10 +84,12 @@ int main() { // LAMBDA: {{%.+}} = alloca i{{[0-9]+}}, // LAMBDA: {{%.+}} = alloca i{{[0-9]+}}, // LAMBDA: [[G_CAST:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[G1_CAST:%.+]] = alloca i{{[0-9]+}}, // LAMBDA: [[SIVAR_CAST:%.+]] = alloca i{{[0-9]+}}, // LAMBDA-DAG: [[G_CAST_VAL:%.+]] = load{{.+}} [[G_CAST]], + // LAMBDA-DAG: [[G1_CAST_VAL:%.+]] = load{{.+}} [[G1_CAST]], // LAMBDA-DAG: [[SIVAR_CAST_VAL:%.+]] = load{{.+}} [[SIVAR_CAST]], - // LAMBDA: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 2, {{.+}} @[[LOUTL1:.+]] to {{.+}}, {{.+}} [[G_CAST_VAL]], {{.+}} [[SIVAR_CAST_VAL]]) + // LAMBDA: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 3, {{.+}} @[[LOUTL1:.+]] to {{.+}}, {{.+}} [[G_CAST_VAL]], {{.+}} [[G1_CAST_VAL]], {{.+}} [[SIVAR_CAST_VAL]]) // LAMBDA: ret void // LAMBDA: define internal void @[[LOUTL1]]({{.+}}) @@ -95,27 +97,24 @@ int main() { // LAMBDA: {{.+}} = alloca i32*, // LAMBDA: {{.+}} = alloca i32*, // LAMBDA: [[G_ADDR:%.+]] = alloca i{{[0-9]+}}, + // LAMBDA: [[G1_ADDR:%.+]] = alloca i{{[0-9]+}}, // LAMBDA: [[SIVAR_ADDR:%.+]] = alloca i{{[0-9]+}}, - // skip loop vars - // LAMBDA: = alloca i32, - // LAMBDA: = alloca i32, - // LAMBDA: = alloca i32, - // LAMBDA: = alloca i32, - // LAMBDA: = alloca i32, - // LAMBDA: [[G1:%.+]] = alloca i{{[0-9]+}}, // LAMBDA: [[G1_TMP:%.+]] = alloca i32*, + // skip loop vars // LAMBDA-DAG: store {{.+}}, {{.+}} [[G_ADDR]], + // LAMBDA-DAG: store {{.+}}, {{.+}} [[G1_ADDR]], // LAMBDA-DAG: store {{.+}}, {{.+}} [[SIVAR_ADDR]], // LAMBDA-DAG: [[G_CONV:%.+]] = bitcast {{.+}} [[G_ADDR]] to + // LAMBDA-DAG: [[G1_CONV:%.+]] = bitcast {{.+}} [[G1_ADDR]] to // LAMBDA-DAG: [[SIVAR_CONV:%.+]] = bitcast {{.+}} [[SIVAR_ADDR]] to - // LAMBDA-DAG: [[G_GBL:%.+]] = load{{.+}}, {{.+}} [[G]], - // LAMBDA-DAG: store{{.+}}, {{.+}} [[G1]], - // LAMBDA-DAG: store{{.+}} [[G1]], {{.+}} [[G1_TMP]], + // LAMBDA-DAG: store{{.+}} [[G1_CONV]], {{.+}} [[G1_TMP]], g = 1; g1 = 1; sivar = 2; // LAMBDA: call void @__kmpc_for_static_init_4( // LAMBDA-DAG: store{{.+}} 1, {{.+}} [[G_CONV]], + // LAMBDA-DAG: [[G1:%.+]] = load{{.+}}, {{.+}}* [[G1_TMP]] + // LAMBDA-DAG: store{{.+}} 1, {{.+}} [[G1]], // LAMBDA-DAG: store{{.+}} 2, {{.+}} [[SIVAR_CONV]], // LAMBDA-DAG: [[G1_REF:%.+]] = load{{.+}}, {{.+}} [[G1_TMP]], // LAMBDA-DAG: store{{.+}} 1, {{.+}} [[G1_REF]], @@ -132,7 +131,10 @@ int main() { // LAMBDA: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // LAMBDA: [[G_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G_PTR_REF]] // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[G_REF]] - // LAMBDA: [[SIVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G1_PTR_REF]] + // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[G1_REF]] + // LAMBDA: [[SIVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // LAMBDA: [[SIVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SIVAR_PTR_REF]] // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SIVAR_REF]] }(); diff --git a/test/OpenMP/teams_distribute_private_codegen.cpp b/test/OpenMP/teams_distribute_private_codegen.cpp index 474c773c2a..06ec35a662 100644 --- a/test/OpenMP/teams_distribute_private_codegen.cpp +++ b/test/OpenMP/teams_distribute_private_codegen.cpp @@ -72,7 +72,7 @@ int main() { // LAMBDA: call void [[OUTER_LAMBDA:@.+]]( [&]() { // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // LAMBDA: call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0)) // LAMBDA: call void @[[LOFFL1:.+]]( // LAMBDA: ret #pragma omp target @@ -118,7 +118,10 @@ int main() { // LAMBDA: [[G_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // LAMBDA: [[G_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G_PTR_REF]] // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[G_REF]] - // LAMBDA: [[SIVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 + // LAMBDA: [[G1_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[G1_PTR_REF]] + // LAMBDA: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[G1_REF]] + // LAMBDA: [[SIVAR_PTR_REF:%.+]] = getelementptr inbounds %{{.+}}, %{{.+}}* [[ARG_PTR]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // LAMBDA: [[SIVAR_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[SIVAR_PTR_REF]] // LAMBDA: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SIVAR_REF]] }(); -- 2.40.0