diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 5aab7722b996..9db8364bb2a8 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -5841,21 +5841,21 @@ emitNumTeamsForTargetDirective(CGOpenMPRuntime &OMPRuntime, const CapturedStmt &CS = *cast(D.getAssociatedStmt()); - // FIXME: Accommodate other combined directives with teams when they become - // available. - if (auto *TeamsDir = dyn_cast_or_null( + if (auto *TeamsDir = dyn_cast_or_null( ignoreCompoundStmts(CS.getCapturedStmt()))) { - if (auto *NTE = TeamsDir->getSingleClause()) { - CGOpenMPInnerExprInfo CGInfo(CGF, CS); - CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo); - llvm::Value *NumTeams = CGF.EmitScalarExpr(NTE->getNumTeams()); - return Bld.CreateIntCast(NumTeams, CGF.Int32Ty, - /*IsSigned=*/true); - } + if (isOpenMPTeamsDirective(TeamsDir->getDirectiveKind())) { + if (auto *NTE = TeamsDir->getSingleClause()) { + CGOpenMPInnerExprInfo CGInfo(CGF, CS); + CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo); + llvm::Value *NumTeams = CGF.EmitScalarExpr(NTE->getNumTeams()); + return Bld.CreateIntCast(NumTeams, CGF.Int32Ty, + /*IsSigned=*/true); + } - // If we have an enclosed teams directive but no num_teams clause we use - // the default value 0. - return Bld.getInt32(0); + // If we have an enclosed teams directive but no num_teams clause we use + // the default value 0. + return Bld.getInt32(0); + } } // No teams associated with the directive. @@ -5944,21 +5944,21 @@ emitNumThreadsForTargetDirective(CGOpenMPRuntime &OMPRuntime, const CapturedStmt &CS = *cast(D.getAssociatedStmt()); - // FIXME: Accommodate other combined directives with teams when they become - // available. - if (auto *TeamsDir = dyn_cast_or_null( + if (auto *TeamsDir = dyn_cast_or_null( ignoreCompoundStmts(CS.getCapturedStmt()))) { - if (auto *TLE = TeamsDir->getSingleClause()) { - CGOpenMPInnerExprInfo CGInfo(CGF, CS); - CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo); - llvm::Value *ThreadLimit = CGF.EmitScalarExpr(TLE->getThreadLimit()); - return CGF.Builder.CreateIntCast(ThreadLimit, CGF.Int32Ty, - /*IsSigned=*/true); - } + if (isOpenMPTeamsDirective(TeamsDir->getDirectiveKind())) { + if (auto *TLE = TeamsDir->getSingleClause()) { + CGOpenMPInnerExprInfo CGInfo(CGF, CS); + CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo); + llvm::Value *ThreadLimit = CGF.EmitScalarExpr(TLE->getThreadLimit()); + return CGF.Builder.CreateIntCast(ThreadLimit, CGF.Int32Ty, + /*IsSigned=*/true); + } - // If we have an enclosed teams directive but no thread_limit clause we use - // the default value 0. - return CGF.Builder.getInt32(0); + // If we have an enclosed teams directive but no thread_limit clause we + // use the default value 0. + return CGF.Builder.getInt32(0); + } } // No teams associated with the directive. diff --git a/clang/test/OpenMP/teams_distribute_codegen.cpp b/clang/test/OpenMP/teams_distribute_codegen.cpp index 53a77de41a1c..8a2afbd4341d 100644 --- a/clang/test/OpenMP/teams_distribute_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_codegen.cpp @@ -25,7 +25,7 @@ int teams_argument_global(int n){ // CK1: [[TE_PAR:%.+]] = load{{.+}}, {{.+}} [[TE_CAST]], // CK1: [[TH_PAR:%.+]] = load{{.+}}, {{.+}} [[TH_CAST]], - // CK1: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // CK1: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 {{.+}}, i32 {{.+}}) // CK1: call void @[[OFFL1:.+]](i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]], #pragma omp target @@ -34,7 +34,7 @@ int teams_argument_global(int n){ a[i] = 0; } - // CK1: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // CK1: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CK1: call void @[[OFFL2:.+]](i{{64|32}} %{{.+}}) #pragma omp target {{{ @@ -93,7 +93,7 @@ int teams_local_arg(void) { int n = 100; int a[n]; - // CK2: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // CK2: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CK2: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) #pragma omp target #pragma omp teams distribute @@ -133,7 +133,7 @@ struct SS{ // CK3: define {{.*}}i32 @{{.+}}foo{{.+}}( int foo(void) { - // CK3: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // CK3: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CK3: call void @[[OFFL1:.+]]([[SSI]]* %{{.+}}) #pragma omp target #pragma omp teams distribute @@ -197,7 +197,7 @@ int main (int argc, char **argv) { } // CK4: define {{.*}}i32 @{{[^,]+}}(i{{.+}}{{.+}} %[[ARGC:.+]], {{.+}}) -// CK4: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0) +// CK4: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CK4: call void @[[OFFL1:.+]]({{.+}}) // CK4: {{%.+}} = call{{.*}} i32 @[[TMAIN:.+]]({{.+}}) // CK4: ret @@ -212,7 +212,7 @@ int main (int argc, char **argv) { // CK4: ret void // CK4: define {{.*}}i32 @[[TMAIN]]({{.+}}) -// CK4: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) +// CK4: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 {{.+}}, i32 {{.+}}) // CK4: call void @[[OFFLT:.+]]({{.+}}) // CK4: ret // CK4-NEXT: } diff --git a/clang/test/OpenMP/teams_distribute_collapse_codegen.cpp b/clang/test/OpenMP/teams_distribute_collapse_codegen.cpp index e9160771e0e1..f4c7a964bc5a 100644 --- a/clang/test/OpenMP/teams_distribute_collapse_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_collapse_codegen.cpp @@ -18,7 +18,7 @@ struct SS{ // CK1: define {{.*}}i32 @{{.+}}foo{{.+}}( int foo(void) { - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL1:.+]]( #pragma omp target #pragma omp teams distribute collapse(2) @@ -89,7 +89,7 @@ int main (int argc, char **argv) { } // CK2: define {{.*}}i32 @{{[^,]+}}(i{{.+}}{{.+}} %[[ARGC:.+]], {{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL1:.+]]({{.+}}) // CK2: {{%.+}} = call{{.*}} i32 @[[TMAIN:.+]]({{.+}}) // CK2: ret @@ -105,7 +105,7 @@ int main (int argc, char **argv) { // CK2: call void @__kmpc_for_static_fini( // CK2: ret void // CK2: define {{.*}}i32 @[[TMAIN]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT1:.+]]({{.+}}) // CK2: ret // CK2-NEXT: } diff --git a/clang/test/OpenMP/teams_distribute_dist_schedule_codegen.cpp b/clang/test/OpenMP/teams_distribute_dist_schedule_codegen.cpp index 685180c26b60..8bd8825e00c4 100644 --- a/clang/test/OpenMP/teams_distribute_dist_schedule_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_dist_schedule_codegen.cpp @@ -18,21 +18,21 @@ struct SS{ // CK1: define {{.*}}i32 @{{.+}}foo{{.+}}( int foo(void) { - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL1:.+]]( #pragma omp target #pragma omp teams distribute for(int i = 0; i < X; i++) { a[i] = (T)0; } - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL2:.+]]( #pragma omp target #pragma omp teams distribute dist_schedule(static) for(int i = 0; i < X; i++) { a[i] = (T)0; } - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL3:.+]]( #pragma omp target #pragma omp teams distribute dist_schedule(static, X/2) @@ -129,11 +129,11 @@ int main (int argc, char **argv) { } // CK2: define {{.*}}i32 @{{[^,]+}}(i{{.+}}{{.+}} %[[ARGC:.+]], {{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL1:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL2:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL3:.+]]({{.+}}) // CK2: {{%.+}} = call{{.*}} i32 @[[TMAIN:.+]]({{.+}}) // CK2: ret @@ -166,11 +166,11 @@ int main (int argc, char **argv) { // CK2: ret void // CK2: define {{.*}}i32 @[[TMAIN]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT1:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT2:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT3:.+]]({{.+}}) // CK2: ret // CK2-NEXT: } diff --git a/clang/test/OpenMP/teams_distribute_firstprivate_codegen.cpp b/clang/test/OpenMP/teams_distribute_firstprivate_codegen.cpp index b9b800e2c0a5..a126b1d16b77 100644 --- a/clang/test/OpenMP/teams_distribute_firstprivate_codegen.cpp +++ b/clang/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(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // LAMBDA: call void @[[LOFFL1:.+]](i{{64|32}} %{{.+}}) // LAMBDA: ret #pragma omp target @@ -154,7 +154,7 @@ int main() { } // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret @@ -248,7 +248,7 @@ int main() { // CHECK: ret void // CHECK: define{{.*}} i{{[0-9]+}} @[[TMAIN_INT]]() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CHECK: call void @[[TOFFL1:.+]](i{{64|32}} %{{.+}}) // CHECK: ret diff --git a/clang/test/OpenMP/teams_distribute_lastprivate_codegen.cpp b/clang/test/OpenMP/teams_distribute_lastprivate_codegen.cpp index cc48ee1109f5..f709b5b7b7fa 100644 --- a/clang/test/OpenMP/teams_distribute_lastprivate_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_lastprivate_codegen.cpp @@ -53,7 +53,7 @@ int main() { [&]() { static float sfvar; // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i{{[0-9]+}} @__tgt_target( + // LAMBDA: call i{{[0-9]+}} @__tgt_target_teams( // LAMBDA: call void [[OFFLOADING_FUN:@.+]]( // LAMBDA: define{{.+}} void [[OFFLOADING_FUN]]( @@ -174,7 +174,7 @@ int main() { // CHECK: define{{.*}} i{{[0-9]+}} @main() // CHECK: [[TEST:%.+]] = alloca [[S_FLOAT_TY]], // CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR:@.+]]([[S_FLOAT_TY]]* [[TEST]]) -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOAD_FUN:@.+]](i{{[0-9]+}} {{.+}}, [2 x i{{[0-9]+}}]* {{.+}}, [2 x [[S_FLOAT_TY]]]* {{.+}}, [[S_FLOAT_TY]]* {{.+}}, i{{[0-9]+}} {{.+}}) // CHECK: ret @@ -271,7 +271,7 @@ int main() { // CHECK: define{{.*}} i{{[0-9]+}} [[TMAIN_INT:@.+]]() // CHECK: [[TEST:%.+]] = alloca [[S_INT_TY]], // CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR:@.+]]([[S_INT_TY]]* [[TEST]]) -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOAD_FUN_1:@.+]](i{{[0-9]+}} {{.+}}, [2 x i{{[0-9]+}}]* {{.+}}, [2 x [[S_INT_TY]]]* {{.+}}, [[S_INT_TY]]* {{.+}}) // CHECK: ret diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp index d2dd9aa8d6cc..d19b832d853e 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp @@ -25,7 +25,7 @@ int teams_argument_global(int n){ // CK1: [[TE_PAR:%.+]] = load{{.+}}, {{.+}} [[TE_CAST]], // CK1: [[TH_PAR:%.+]] = load{{.+}}, {{.+}} [[TH_CAST]], - // CK1: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // CK1: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 {{.+}}, i32 {{.+}}) // CK1: call void @[[OFFL1:.+]](i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]], #pragma omp target @@ -35,7 +35,7 @@ int teams_argument_global(int n){ #pragma omp cancel for } - // CK1: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // CK1: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CK1: call void @[[OFFL2:.+]](i{{64|32}} %{{.+}}) #pragma omp target {{{ @@ -96,7 +96,7 @@ int teams_local_arg(void) { int n = 100; int a[n]; - // CK2: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // CK2: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CK2: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) #pragma omp target #pragma omp teams distribute parallel for @@ -137,7 +137,7 @@ struct SS{ // CK3: define {{.*}}i32 @{{.+}}foo{{.+}}( int foo(void) { - // CK3: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // CK3: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CK3: call void @[[OFFL1:.+]]([[SSI]]* %{{.+}}) #pragma omp target #pragma omp teams distribute parallel for @@ -202,7 +202,7 @@ int main (int argc, char **argv) { } // CK4: define {{.*}}i32 @{{[^,]+}}(i{{.+}}{{.+}} %[[ARGC:.+]], {{.+}}) -// CK4: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0) +// CK4: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CK4: call void @[[OFFL1:.+]]({{.+}}) // CK4: {{%.+}} = call{{.*}} i32 @[[TMAIN:.+]]({{.+}}) // CK4: ret @@ -218,7 +218,7 @@ int main (int argc, char **argv) { // CK4: ret void // CK4: define {{.*}}i32 @[[TMAIN]]({{.+}}) -// CK4: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) +// CK4: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 {{.+}}, i32 {{.+}}) // CK4: call void @[[OFFLT:.+]]({{.+}}) // CK4: ret // CK4-NEXT: } diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_collapse_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_collapse_codegen.cpp index 3aad95a91d81..7522288f75e2 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_collapse_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_collapse_codegen.cpp @@ -18,13 +18,13 @@ struct SS{ // CK1: define {{.*}}i32 @{{.+}}foo{{.+}}( int foo(void) { - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL1:.+]]( #pragma omp target #pragma omp teams distribute parallel for collapse(2) for(int i = 0; i < X; i++) { for(int j = 0; j < Y; j++) { - a[i][j] = (T)0; + a[i][j] = (T)0; } } // CK1: define internal void @[[OFFL1]]( @@ -93,7 +93,7 @@ int main (int argc, char **argv) { } // CK2: define {{.*}}i32 @{{[^,]+}}(i{{.+}}{{.+}} %[[ARGC:.+]], {{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL1:.+]]({{.+}}) // CK2: {{%.+}} = call{{.*}} i32 @[[TMAIN:.+]]({{.+}}) // CK2: ret @@ -117,7 +117,7 @@ int main (int argc, char **argv) { // CK2: define {{.*}}i32 @[[TMAIN]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT1:.+]]({{.+}}) // CK2: ret // CK2-NEXT: } diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_copyin_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_copyin_codegen.cpp index 6d0f7af7bbe5..0202a40002b5 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_copyin_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_copyin_codegen.cpp @@ -34,7 +34,7 @@ int main() { // LAMBDA: call void [[OUTER_LAMBDA:@.+]]( [&]() { // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i32 @__tgt_target( + // LAMBDA: call i32 @__tgt_target_teams( // LAMBDA: call void @[[LOFFL1:.+]]( // LAMBDA: ret #pragma omp target @@ -102,7 +102,7 @@ int main() { } // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target( +// CHECK: call i32 @__tgt_target_teams( // CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret @@ -148,7 +148,7 @@ int main() { // CHECK: ret void // CHECK: define{{.*}} i{{[0-9]+}} @[[TMAIN_INT]]() -// CHECK: call i32 @__tgt_target( +// CHECK: call i32 @__tgt_target_teams( // CHECK: call void @[[TOFFL1:.+]]( // CHECK: ret diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_dist_schedule_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_dist_schedule_codegen.cpp index 717effd89bf2..dc48f166b4c0 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_dist_schedule_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_dist_schedule_codegen.cpp @@ -18,21 +18,21 @@ struct SS{ // CK1: define {{.*}}i32 @{{.+}}foo{{.+}}( int foo(void) { - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL1:.+]]( #pragma omp target #pragma omp teams distribute parallel for for(int i = 0; i < X; i++) { a[i] = (T)0; } - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL2:.+]]( #pragma omp target #pragma omp teams distribute parallel for dist_schedule(static) for(int i = 0; i < X; i++) { a[i] = (T)0; } - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL3:.+]]( #pragma omp target #pragma omp teams distribute parallel for dist_schedule(static, X/2) @@ -150,11 +150,11 @@ int main (int argc, char **argv) { } // CK2: define {{.*}}i32 @{{[^,]+}}(i{{.+}}{{.+}} %[[ARGC:.+]], {{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL1:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL2:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL3:.+]]({{.+}}) // CK2: {{%.+}} = call{{.*}} i32 @[[TMAIN:.+]]({{.+}}) // CK2: ret @@ -205,11 +205,11 @@ int main (int argc, char **argv) { // CK2: ret void // CK2: define {{.*}}i32 @[[TMAIN]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT1:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT2:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT3:.+]]({{.+}}) // CK2: ret // CK2-NEXT: } diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_firstprivate_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_firstprivate_codegen.cpp index 6ee5bfc6b4b6..55ef18fc6bcc 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_firstprivate_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_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(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // LAMBDA: call void @[[LOFFL1:.+]](i{{64|32}} %{{.+}}) // LAMBDA: ret #pragma omp target @@ -187,7 +187,7 @@ int main() { } // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret @@ -347,7 +347,7 @@ int main() { // CHECK: ret void // CHECK: define{{.*}} i{{[0-9]+}} @[[TMAIN_INT]]() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CHECK: call void @[[TOFFL1:.+]](i{{64|32}} %{{.+}}) // CHECK: ret diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_if_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_if_codegen.cpp index 495434cc9c9a..bc40dca83fbe 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_if_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_if_codegen.cpp @@ -17,9 +17,9 @@ int Arg; // CHECK-LABEL: define {{.*}}void @{{.+}}gtid_test void gtid_test() { #pragma omp target -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOADING_FUN_0:@.+]]( -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOADING_FUN_1:@.+]]( #pragma omp teams distribute parallel for for(int i = 0 ; i < 100; i++) {} @@ -77,11 +77,11 @@ int tmain(T Arg) { // CHECK-LABEL: define {{.*}}i{{[0-9]+}} @main() int main() { -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOADING_FUN_0:@.+]]( -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOADING_FUN_1:@.+]]( -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOADING_FUN_2:@.+]]( // CHECK: = call {{.*}}i{{.+}} @{{.+}}tmain #pragma omp target diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_lastprivate_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_lastprivate_codegen.cpp index 454e80c257b5..57e98f1b5c6e 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_lastprivate_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_lastprivate_codegen.cpp @@ -53,7 +53,7 @@ int main() { [&]() { static float sfvar; // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i{{[0-9]+}} @__tgt_target( + // LAMBDA: call i{{[0-9]+}} @__tgt_target_teams( // LAMBDA: call void [[OFFLOADING_FUN:@.+]]( // LAMBDA: define{{.+}} void [[OFFLOADING_FUN]]( @@ -227,7 +227,7 @@ int main() { // CHECK: define{{.*}} i{{[0-9]+}} @main() // CHECK: [[TEST:%.+]] = alloca [[S_FLOAT_TY]], // CHECK: call {{.*}} [[S_FLOAT_TY_DEF_CONSTR:@.+]]([[S_FLOAT_TY]]* [[TEST]]) -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOAD_FUN:@.+]](i{{[0-9]+}} {{.+}}, [2 x i{{[0-9]+}}]* {{.+}}, [2 x [[S_FLOAT_TY]]]* {{.+}}, [[S_FLOAT_TY]]* {{.+}}, i{{[0-9]+}} {{.+}}) // CHECK: ret @@ -404,7 +404,7 @@ int main() { // CHECK: define{{.*}} i{{[0-9]+}} [[TMAIN_INT:@.+]]() // CHECK: [[TEST:%.+]] = alloca [[S_INT_TY]], // CHECK: call {{.*}} [[S_INT_TY_DEF_CONSTR:@.+]]([[S_INT_TY]]* [[TEST]]) -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOAD_FUN_1:@.+]](i{{[0-9]+}} {{.+}}, [2 x i{{[0-9]+}}]* {{.+}}, [2 x [[S_INT_TY]]]* {{.+}}, [[S_INT_TY]]* {{.+}}) // CHECK: ret diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_num_threads_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_num_threads_codegen.cpp index 1ca7db895378..7c0ac461220d 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_num_threads_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_num_threads_codegen.cpp @@ -37,9 +37,9 @@ int tmain() { int main() { S s(0); char a = s; -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOADING_FUN_0:@.+]]( -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[OFFLOADING_FUN_1:@.+]]( // CHECK: invoke{{.+}} [[TMAIN_5:@.+]]() // CHECK: invoke{{.+}} [[TMAIN_1:@.+]]() @@ -73,16 +73,16 @@ int main() { // tmain 5 // CHECK-DAG: define {{.*}}i{{[0-9]+}} [[TMAIN_5]]() -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[T_OFFLOADING_FUN_0:@.+]]( -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[T_OFFLOADING_FUN_1:@.+]]( // tmain 1 // CHECK-DAG: define {{.*}}i{{[0-9]+}} [[TMAIN_1]]() -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[T_OFFLOADING_FUN_2:@.+]]( -// CHECK: call i{{[0-9]+}} @__tgt_target( +// CHECK: call i{{[0-9]+}} @__tgt_target_teams( // CHECK: call void [[T_OFFLOADING_FUN_3:@.+]]( // CHECK: define internal void [[T_OFFLOADING_FUN_0]]( diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_private_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_private_codegen.cpp index 5a5508458115..97541315bff5 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_private_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_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(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // LAMBDA: call void @[[LOFFL1:.+]]( // LAMBDA: ret #pragma omp target @@ -162,7 +162,7 @@ int main() { } // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret @@ -238,7 +238,7 @@ int main() { // CHECK: ret void // CHECK: define{{.*}} i{{[0-9]+}} @[[TMAIN_INT]]() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 0, +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, // CHECK: call void @[[TOFFL1:.+]]() // CHECK: ret diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_proc_bind_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_proc_bind_codegen.cpp index ed7f2dd39a7f..71e92bd5a555 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_proc_bind_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_proc_bind_codegen.cpp @@ -41,9 +41,9 @@ int main() { return tmain(); } -// CHECK: call {{.*}}@__tgt_target({{.+}}) +// CHECK: call {{.*}}@__tgt_target_teams({{.+}}) // CHECK: call void [[OFFL1:@.+]]() -// CHECK: call {{.*}}@__tgt_target({{.+}}) +// CHECK: call {{.*}}@__tgt_target_teams({{.+}}) // CHECK: call void [[OFFL2:@.+]]() // CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]() // CHECK: ret i32 [[CALL_RET]] @@ -73,7 +73,7 @@ int main() { // CHECK: ret void // CHECK: define{{.+}} [[TMAIN]]() -// CHECK: call {{.*}}@__tgt_target({{.+}}) +// CHECK: call {{.*}}@__tgt_target_teams({{.+}}) // CHECK: call void [[OFFL3:@.+]]() // CHECK: define{{.+}} [[OFFL3]]() diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_reduction_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_reduction_codegen.cpp index e33f0c5263c8..09dc3ec8e17c 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_reduction_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_reduction_codegen.cpp @@ -34,7 +34,7 @@ int main() { // LAMBDA: call void [[OUTER_LAMBDA:@.+]]( [&]() { // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // LAMBDA: call void @[[LOFFL1:.+]]( // LAMBDA: ret #pragma omp target @@ -158,7 +158,7 @@ int main() { // CHECK: [[RED_VAR:@.+]] = common global [8 x {{.+}}] zeroinitializer // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret @@ -250,7 +250,7 @@ int main() { // CHECK: br // CHECK: define{{.*}} i{{[0-9]+}} @[[TMAIN_INT]]() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 1, +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, // CHECK: call void @[[TOFFL1:.+]]({{.+}}) // CHECK: ret diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_schedule_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_schedule_codegen.cpp index f5a00cd30e0c..201a170ccd4a 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_schedule_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_schedule_codegen.cpp @@ -18,21 +18,21 @@ struct SS{ // CK1: define {{.*}}i32 @{{.+}}foo{{.+}}( int foo(void) { - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL1:.+]]( #pragma omp target #pragma omp teams distribute parallel for for(int i = 0; i < X; i++) { a[i] = (T)0; } - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL2:.+]]( #pragma omp target #pragma omp teams distribute parallel for schedule(static) for(int i = 0; i < X; i++) { a[i] = (T)0; } - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL3:.+]]( #pragma omp target #pragma omp teams distribute parallel for schedule(static, X/2) @@ -40,7 +40,7 @@ struct SS{ a[i] = (T)0; } - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL4:.+]]( #pragma omp target #pragma omp teams distribute parallel for schedule(dynamic) @@ -48,7 +48,7 @@ struct SS{ a[i] = (T)0; } - // CK1: call i32 @__tgt_target( + // CK1: call i32 @__tgt_target_teams( // CK1: call void @[[OFFL5:.+]]( #pragma omp target #pragma omp teams distribute parallel for schedule(dynamic, X/2) @@ -216,15 +216,15 @@ int main (int argc, char **argv) { } // CK2: define {{.*}}i32 @{{[^,]+}}(i{{.+}}{{.+}} %[[ARGC:.+]], {{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL1:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL2:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL3:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL4:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFL5:.+]]({{.+}}) // CK2: {{%.+}} = call{{.*}} i32 @[[TMAIN:.+]]({{.+}}) // CK2: ret @@ -306,15 +306,15 @@ int main (int argc, char **argv) { // CK2: ret void // CK2: define {{.*}}i32 @[[TMAIN]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT1:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT2:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT3:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT4:.+]]({{.+}}) -// CK2: call i32 @__tgt_target( +// CK2: call i32 @__tgt_target_teams( // CK2: call void @[[OFFLT5:.+]]({{.+}}) // CK2: ret // CK2-NEXT: } diff --git a/clang/test/OpenMP/teams_distribute_private_codegen.cpp b/clang/test/OpenMP/teams_distribute_private_codegen.cpp index e2dfb5de506e..23d1b9067e1b 100644 --- a/clang/test/OpenMP/teams_distribute_private_codegen.cpp +++ b/clang/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(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // LAMBDA: call void @[[LOFFL1:.+]]( // LAMBDA: ret #pragma omp target @@ -141,7 +141,7 @@ int main() { } // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret @@ -186,7 +186,7 @@ int main() { // CHECK: define{{.*}} i{{[0-9]+}} @[[TMAIN_INT]]() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 0, +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 0, // CHECK: call void @[[TOFFL1:.+]]() // CHECK: ret diff --git a/clang/test/OpenMP/teams_distribute_reduction_codegen.cpp b/clang/test/OpenMP/teams_distribute_reduction_codegen.cpp index 2360e1be8bcc..961891da98d0 100644 --- a/clang/test/OpenMP/teams_distribute_reduction_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_reduction_codegen.cpp @@ -34,7 +34,7 @@ int main() { // LAMBDA: call void [[OUTER_LAMBDA:@.+]]( [&]() { // LAMBDA: define{{.*}} internal{{.*}} void [[OUTER_LAMBDA]]( - // LAMBDA: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) + // LAMBDA: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // LAMBDA: call void @[[LOFFL1:.+]]( // LAMBDA: ret #pragma omp target @@ -115,7 +115,7 @@ int main() { // CHECK: [[RED_VAR:@.+]] = common global [8 x {{.+}}] zeroinitializer // CHECK: define {{.*}}i{{[0-9]+}} @main() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0)) +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) // CHECK: call void @[[OFFL1:.+]](i{{64|32}} %{{.+}}) // CHECK: {{%.+}} = call{{.*}} i32 @[[TMAIN_INT:.+]]() // CHECK: ret @@ -166,7 +166,7 @@ int main() { // CHECK: define{{.*}} i{{[0-9]+}} @[[TMAIN_INT]]() -// CHECK: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 1, +// CHECK: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, // CHECK: call void @[[TOFFL1:.+]]({{.+}}) // CHECK: ret