Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -113,17 +113,21 @@ bool CancelRegion; unsigned AssociatedLoops; SourceLocation InnerTeamsRegionLoc; + const DeclContext *ParentDeclContext; SharingMapTy(OpenMPDirectiveKind DKind, DeclarationNameInfo Name, - Scope *CurScope, SourceLocation Loc) + Scope *CurScope, SourceLocation Loc, + const DeclContext *ParentDeclContext) : SharingMap(), AlignedMap(), LCVMap(), DefaultAttr(DSA_unspecified), Directive(DKind), DirectiveName(std::move(Name)), CurScope(CurScope), ConstructLoc(Loc), OrderedRegion(), NowaitRegion(false), - CancelRegion(false), AssociatedLoops(1), InnerTeamsRegionLoc() {} + CancelRegion(false), AssociatedLoops(1), InnerTeamsRegionLoc(), + ParentDeclContext(ParentDeclContext) {} SharingMapTy() : SharingMap(), AlignedMap(), LCVMap(), DefaultAttr(DSA_unspecified), Directive(OMPD_unknown), DirectiveName(), CurScope(nullptr), ConstructLoc(), OrderedRegion(), NowaitRegion(false), - CancelRegion(false), AssociatedLoops(1), InnerTeamsRegionLoc() {} + CancelRegion(false), AssociatedLoops(1), InnerTeamsRegionLoc(), + ParentDeclContext(nullptr) {} }; typedef SmallVector StackTy; @@ -156,8 +160,10 @@ void setForceVarCapturing(bool V) { ForceCapturing = V; } void push(OpenMPDirectiveKind DKind, const DeclarationNameInfo &DirName, - Scope *CurScope, SourceLocation Loc) { - Stack.push_back(SharingMapTy(DKind, DirName, CurScope, Loc)); + Scope *CurScope, SourceLocation Loc, + const DeclContext *ParentDeclContext) { + Stack.push_back( + SharingMapTy(DKind, DirName, CurScope, Loc, ParentDeclContext)); Stack.back().DefaultAttrLoc = Loc; } @@ -245,8 +251,9 @@ return Stack[Stack.size() - 2].Directive; return OMPD_unknown; } - /// \brief Return the directive associated with the provided scope. - OpenMPDirectiveKind getDirectiveForScope(const Scope *S) const; + /// \brief Return the directive associated with the provided region scope. + OpenMPDirectiveKind + getDirectiveForRegionScope(const CapturedRegionScopeInfo *RSI) const; /// \brief Set default data sharing attribute to none. void setDefaultDSANone(SourceLocation Loc) { @@ -785,6 +792,9 @@ template bool DSAStackTy::hasDirective(NamedDirectivesPredicate DPred, bool FromParent) { + // We look only in the enclosing region. + if (Stack.size() < 2) + return false; auto StartI = std::next(Stack.rbegin()); auto EndI = std::prev(Stack.rend()); if (FromParent && StartI != EndI) { @@ -797,9 +807,16 @@ return false; } -OpenMPDirectiveKind DSAStackTy::getDirectiveForScope(const Scope *S) const { +OpenMPDirectiveKind DSAStackTy::getDirectiveForRegionScope( + const CapturedRegionScopeInfo *RSI) const { + // Directives are expected to have a context associated. + if (!(RSI && RSI->TheCapturedDecl)) + return OMPD_unknown; + + // A DSA refers to this captured region if the parent contexts match. + auto *ParentContext = RSI->TheCapturedDecl->getParent(); for (auto I = Stack.rbegin(), EE = Stack.rend(); I != EE; ++I) - if (I->CurScope == S) + if (I->ParentDeclContext == ParentContext) return I->Directive; return OMPD_unknown; } @@ -818,7 +835,7 @@ bool IsByRef = true; // Find the directive that is associated with the provided scope. - auto DKind = DSAStack->getDirectiveForScope(RSI->TheScope); + auto DKind = DSAStack->getDirectiveForRegionScope(RSI); auto Ty = D->getType(); if (isOpenMPTargetExecutionDirective(DKind)) { @@ -914,8 +931,7 @@ if (DSAStack->getCurrentDirective() == OMPD_target && !DSAStack->isClauseParsingMode()) return VD; - if (DSAStack->getCurScope() && - DSAStack->hasDirective( + if (DSAStack->hasDirective( [](OpenMPDirectiveKind K, const DeclarationNameInfo &DNI, SourceLocation Loc) -> bool { return isOpenMPTargetExecutionDirective(K); @@ -964,7 +980,9 @@ void Sema::StartOpenMPDSABlock(OpenMPDirectiveKind DKind, const DeclarationNameInfo &DirName, Scope *CurScope, SourceLocation Loc) { - DSAStack->push(DKind, DirName, CurScope, Loc); + // The current Sema declarative context is the parent of the captured region + // that refers to this DSA block. + DSAStack->push(DKind, DirName, CurScope, Loc, CurContext); PushExpressionEvaluationContext(PotentiallyEvaluated); } Index: test/OpenMP/target_codegen_global_capture.cpp =================================================================== --- test/OpenMP/target_codegen_global_capture.cpp +++ test/OpenMP/target_codegen_global_capture.cpp @@ -21,6 +21,11 @@ // CHECK-DAG: [[BB:@.+]] = internal global float 1.000000e+01 // CHECK-DAG: [[BC:@.+]] = internal global float 1.100000e+01 // CHECK-DAG: [[BD:@.+]] = internal global float 1.200000e+01 +// CHECK-DAG: [[TBA:@.+]] = {{.*}}global float 1.700000e+01 +// CHECK-DAG: [[TBB:@.+]] = {{.*}}global float 1.800000e+01 +// CHECK-DAG: [[TBC:@.+]] = {{.*}}global float 1.900000e+01 +// CHECK-DAG: [[TBD:@.+]] = {{.*}}global float 2.000000e+01 + double Ga = 1.0; double Gb = 2.0; double Gc = 3.0; @@ -42,14 +47,14 @@ static float Sd = 8.0; // CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LB]], - // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* @Gb, - // CHECK-DAG: [[VALFB:%.+]] = load float, float* @_ZZ3foossssE2Sb, - // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* @Gc, + // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]], + // CHECK-DAG: [[VALFB:%.+]] = load float, float* [[FB]], + // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]], // CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LC]], - // CHECK-DAG: [[VALFC:%.+]] = load float, float* @_ZZ3foossssE2Sc, + // CHECK-DAG: [[VALFC:%.+]] = load float, float* [[FC]], // CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LD]], - // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* @Gd, - // CHECK-DAG: [[VALFD:%.+]] = load float, float* @_ZZ3foossssE2Sd, + // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]], + // CHECK-DAG: [[VALFD:%.+]] = load float, float* [[FD]], // 3 local vars being captured. @@ -178,14 +183,156 @@ #pragma omp parallel { // CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LLB]], - // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* @Gb, - // CHECK-DAG: [[VALFB:%.+]] = load float, float* @_ZZ3barssssE2Sb, - // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* @Gc, + // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]], + // CHECK-DAG: [[VALFB:%.+]] = load float, float* [[BB]], + // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]], + // CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LLC]], + // CHECK-DAG: [[VALFC:%.+]] = load float, float* [[BC]], + // CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LLD]], + // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]], + // CHECK-DAG: [[VALFD:%.+]] = load float, float* [[BD]], + + // 3 local vars being captured. + + // CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]], + // CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16* + // CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]], + // CHECK-DAG: [[CPTRLB:%.+]] = inttoptr i[[sz]] [[CVALLB]] to i8* + // CHECK-DAG: store i8* [[CPTRLB]], i8** [[GEPLB:%.+]], + // CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}} + + // CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]], + // CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16* + // CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]], + // CHECK-DAG: [[CPTRLC:%.+]] = inttoptr i[[sz]] [[CVALLC]] to i8* + // CHECK-DAG: store i8* [[CPTRLC]], i8** [[GEPLC:%.+]], + // CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}} + + // CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]], + // CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16* + // CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]], + // CHECK-DAG: [[CPTRLD:%.+]] = inttoptr i[[sz]] [[CVALLD]] to i8* + // CHECK-DAG: store i8* [[CPTRLD]], i8** [[GEPLD:%.+]], + // CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}} + + // 3 static vars being captured. + + // CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]], + // CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float* + // CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]], + // CHECK-DAG: [[CPTRFB:%.+]] = inttoptr i[[sz]] [[CVALFB]] to i8* + // CHECK-DAG: store i8* [[CPTRFB]], i8** [[GEPFB:%.+]], + // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}} + + // CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]], + // CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float* + // CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]], + // CHECK-DAG: [[CPTRFC:%.+]] = inttoptr i[[sz]] [[CVALFC]] to i8* + // CHECK-DAG: store i8* [[CPTRFC]], i8** [[GEPFC:%.+]], + // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}} + + // CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]], + // CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float* + // CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]], + // CHECK-DAG: [[CPTRFD:%.+]] = inttoptr i[[sz]] [[CVALFD]] to i8* + // CHECK-DAG: store i8* [[CPTRFD]], i8** [[GEPFD:%.+]], + // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}} + + // 3 static global vars being captured. + + // CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]], + // CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double* + // CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]], + // CHECK-64-DAG: [[CPTRGB:%.+]] = inttoptr i[[sz]] [[CVALGB]] to i8* + // CHECK-64-DAG: store i8* [[CPTRGB]], i8** [[GEPGB:%.+]], + // CHECK-32-DAG: store i8* bitcast (double* @Gb to i8*), i8** [[GEPGB:%.+]], + // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}} + + // CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]], + // CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double* + // CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]], + // CHECK-64-DAG: [[CPTRGC:%.+]] = inttoptr i[[sz]] [[CVALGC]] to i8* + // CHECK-64-DAG: store i8* [[CPTRGC]], i8** [[GEPGC:%.+]], + // CHECK-32-DAG: store i8* bitcast (double* @Gc to i8*), i8** [[GEPGC:%.+]], + // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}} + + // CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]], + // CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double* + // CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]], + // CHECK-64-DAG: [[CPTRGD:%.+]] = inttoptr i[[sz]] [[CVALGD]] to i8* + // CHECK-64-DAG: store i8* [[CPTRGD]], i8** [[GEPGD:%.+]], + // CHECK-32-DAG: store i8* bitcast (double* @Gd to i8*), i8** [[GEPGD:%.+]], + // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}} + + // CHECK: call i32 @__tgt_target + // CHECK: call void [[OFFLOADF:@.+]]( + // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd + #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0) + { + b += 1; + Gb += 1.0; + Sb += 1.0; + + // CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}) + // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}) + + // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}) + // Capture d, Gd, Sd + #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0) + { + d += 1; + Gd += 1.0; + Sd += 1.0; + } + } + } + return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd; +} + +/// +/// Tests with template functions. +/// + +// CHECK: define {{.*}} @{{.*}}tbar2{{.*}}( + +// CHECK: define {{.*}} @{{.*}}tbar{{.*}}( +// CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]], +// CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]], +// CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]], +// CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]]) +// CHECK: [[LA:%.+]] = alloca i16 +// CHECK: [[LB:%.+]] = alloca i16 +// CHECK: [[LC:%.+]] = alloca i16 +// CHECK: [[LD:%.+]] = alloca i16 +template +int tbar(T a, T b, T c, T d){ + static float Sa = 17.0; + static float Sb = 18.0; + static float Sc = 19.0; + static float Sd = 20.0; + + // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}}) + // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, i16* dereferenceable(2) [[A:%.+]], i16* dereferenceable(2) [[B:%.+]], i16* dereferenceable(2) [[C:%.+]], i16* dereferenceable(2) [[D:%.+]]) + // Capture a, b, c, d + // CHECK: [[ALLOCLA:%.+]] = alloca i16 + // CHECK: [[ALLOCLB:%.+]] = alloca i16 + // CHECK: [[ALLOCLC:%.+]] = alloca i16 + // CHECK: [[ALLOCLD:%.+]] = alloca i16 + // CHECK: [[LLA:%.+]] = load i16*, i16** [[ALLOCLA]], + // CHECK: [[LLB:%.+]] = load i16*, i16** [[ALLOCLB]], + // CHECK: [[LLC:%.+]] = load i16*, i16** [[ALLOCLC]], + // CHECK: [[LLD:%.+]] = load i16*, i16** [[ALLOCLD]], + #pragma omp parallel + { + // CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LLB]], + // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]], + // CHECK-DAG: [[VALFB:%.+]] = load float, float* [[TBB]], + // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]], // CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LLC]], - // CHECK-DAG: [[VALFC:%.+]] = load float, float* @_ZZ3barssssE2Sc, + // CHECK-DAG: [[VALFC:%.+]] = load float, float* [[TBC]], // CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LLD]], - // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* @Gd, - // CHECK-DAG: [[VALFD:%.+]] = load float, float* @_ZZ3barssssE2Sd, + // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]], + // CHECK-DAG: [[VALFD:%.+]] = load float, float* [[TBD]], // 3 local vars being captured. @@ -284,4 +431,8 @@ return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd; } +int tbar2(short a, short b, short c, short d){ + return tbar(a, b, c, d); +} + #endif Index: test/OpenMP/teams_codegen.cpp =================================================================== --- test/OpenMP/teams_codegen.cpp +++ test/OpenMP/teams_codegen.cpp @@ -250,12 +250,10 @@ // CK4: ret void // CK4-NEXT: } -// CK4: define {{.*}}void @{{[^,]+}}(i8*** dereferenceable({{.}}) [[ARGC1:%.+]]) -// CK4: [[ARGCADDR1:%.+]] = alloca i8*** -// CK4: store i8*** [[ARGC1]], i8**** [[ARGCADDR1]] -// CK4: [[CONV1:%.+]] = load i8***, i8**** [[ARGCADDR1]] -// CK4: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* {{.+}} to void (i32*, i32*, ...)*), i8*** [[CONV1]]) - +// CK4: define {{.*}}void @{{[^,]+}}(i8** [[ARGC1:%.+]]) +// CK4: [[ARGCADDR1:%.+]] = alloca i8** +// CK4: store i8** [[ARGC1]], i8*** [[ARGCADDR1]] +// CK4: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* {{.+}} to void (i32*, i32*, ...)*), i8*** [[ARGCADDR1]]) #endif // CK4 @@ -320,21 +318,22 @@ // CK5-64: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[CONV]]) // CK5-32: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[ARGCADDR]]) -// CK5: define {{.*}}void @{{[^,]+}}(i{{.+}} dereferenceable({{.+}}) [[AP:%.+]], i{{.+}} dereferenceable({{.+}}) [[BP:%.+]], i{{.+}} dereferenceable({{.+}}) [[ARGC:%.+]]) +// CK5: define {{.*}}void @{{[^,]+}}(i{{.+}} [[AP:%.+]], i{{.+}} [[BP:%.+]], i8** [[ARGC:%.+]]) // CK5: [[AADDR:%.+]] = alloca i{{.+}} // CK5: [[BADDR:%.+]] = alloca i{{.+}} -// CK5: [[ARGCADDR:%.+]] = alloca i{{.+}} +// CK5: [[ARGCADDR:%.+]] = alloca i8** // CK5: [[GBL_TH_NUM:%.+]] = call i32 @__kmpc_global_thread_num(%ident_t* [[DEF_LOC_0]]) // CK5: store i{{.+}} [[AP]], i{{.+}}* [[AADDR]] // CK5: store i{{.+}} [[BP]], i{{.+}}* [[BADDR]] -// CK5: store i{{.+}} [[ARGC]], i{{.+}}* [[ARGCADDR]] -// CK5: [[A_ADDR_VAL:%.+]] = load i32*, i32** [[AADDR]] -// CK5: [[B_ADDR_VAL:%.+]] = load i32*, i32** [[BADDR]] -// CK5: [[ARGC_ADDR_VAL:%.+]] = load i{{.+}}, i{{.+}}* [[ARGCADDR]] -// CK5: [[A_VAL:%.+]] = load i32, i32* [[A_ADDR_VAL]] -// CK5: [[B_VAL:%.+]] = load i32, i32* [[B_ADDR_VAL]] -// CK5: {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[A_VAL]], i32 [[B_VAL]]) -// CK5: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i{{.+}})* @.omp_outlined.{{.+}} to void (i32*, i32*, ...)*), i{{.+}} [[ARGC_ADDR_VAL]]) +// CK5: store i8** [[ARGC]], i8*** [[ARGCADDR]] +// CK5-64: [[ACONV:%.+]] = bitcast i64* [[AADDR]] to i32* +// CK5-64: [[BCONV:%.+]] = bitcast i64* [[BADDR]] to i32* +// CK5-64: [[ACONVVAL:%.+]] = load i32, i32* [[ACONV]] +// CK5-64: [[BCONVVAL:%.+]] = load i32, i32* [[BCONV]] +// CK5-32: [[ACONVVAL:%.+]] = load i32, i32* [[AADDR]] +// CK5-32: [[BCONVVAL:%.+]] = load i32, i32* [[BADDR]] +// CK5: {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[ACONVVAL]], i32 [[BCONVVAL]]) +// CK5: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i{{.+}})* @.omp_outlined.{{.+}} to void (i32*, i32*, ...)*), i8*** [[ARGCADDR]]) // CK5: ret void // CK5-NEXT: }