replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[OPENMP] Emit `__tgt_target_teams` for all teams directives. 

Previously we emitted `__tgt_target_teams` only for standalone teams
directives. This patch allows emit this function for all teams-based
directives.

llvm-svn: 319585
This commit is contained in:
Alexey Bataev 2017-12-01 21:31:08 +00:00
parent 7bfcc05830
commit 50a1c7860f
20 changed files with 119 additions and 119 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

52
clang/lib/CodeGen/CGOpenMPRuntime.cpp
View File

@ -5841,21 +5841,21 @@ emitNumTeamsForTargetDirective(CGOpenMPRuntime &OMPRuntime,
const CapturedStmt &CS = *cast<CapturedStmt>(D.getAssociatedStmt());
// FIXME: Accommodate other combined directives with teams when they become
// available.
if (auto *TeamsDir = dyn_cast_or_null<OMPTeamsDirective>(
if (auto *TeamsDir = dyn_cast_or_null<OMPExecutableDirective>(
ignoreCompoundStmts(CS.getCapturedStmt()))) {
if (auto *NTE = TeamsDir->getSingleClause<OMPNumTeamsClause>()) {
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<OMPNumTeamsClause>()) {
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<CapturedStmt>(D.getAssociatedStmt());
// FIXME: Accommodate other combined directives with teams when they become
// available.
if (auto *TeamsDir = dyn_cast_or_null<OMPTeamsDirective>(
if (auto *TeamsDir = dyn_cast_or_null<OMPExecutableDirective>(
ignoreCompoundStmts(CS.getCapturedStmt()))) {
if (auto *TLE = TeamsDir->getSingleClause<OMPThreadLimitClause>()) {
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<OMPThreadLimitClause>()) {
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.

12
clang/test/OpenMP/teams_distribute_codegen.cpp
View File

@ -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: }

6
clang/test/OpenMP/teams_distribute_collapse_codegen.cpp
View File

@ -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: }

18
clang/test/OpenMP/teams_distribute_dist_schedule_codegen.cpp
View File

@ -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: }

6
clang/test/OpenMP/teams_distribute_firstprivate_codegen.cpp
View File

@ -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

6
clang/test/OpenMP/teams_distribute_lastprivate_codegen.cpp
View File

@ -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

12
clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp
View File

@ -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: }

8
clang/test/OpenMP/teams_distribute_parallel_for_collapse_codegen.cpp
View File

@ -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: }

6
clang/test/OpenMP/teams_distribute_parallel_for_copyin_codegen.cpp
View File

@ -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

18
clang/test/OpenMP/teams_distribute_parallel_for_dist_schedule_codegen.cpp
View File

@ -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: }

6
clang/test/OpenMP/teams_distribute_parallel_for_firstprivate_codegen.cpp
View File

@ -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

10
clang/test/OpenMP/teams_distribute_parallel_for_if_codegen.cpp
View File

@ -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

6
clang/test/OpenMP/teams_distribute_parallel_for_lastprivate_codegen.cpp
View File

@ -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

12
clang/test/OpenMP/teams_distribute_parallel_for_num_threads_codegen.cpp
View File

@ -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]](

6
clang/test/OpenMP/teams_distribute_parallel_for_private_codegen.cpp
View File

@ -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

6
clang/test/OpenMP/teams_distribute_parallel_for_proc_bind_codegen.cpp
View File

@ -41,9 +41,9 @@ int main() {
return tmain<int>();
}
// 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]]()

6
clang/test/OpenMP/teams_distribute_parallel_for_reduction_codegen.cpp
View File

@ -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

30
clang/test/OpenMP/teams_distribute_parallel_for_schedule_codegen.cpp
View File

@ -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: }

6
clang/test/OpenMP/teams_distribute_private_codegen.cpp
View File

@ -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

6
clang/test/OpenMP/teams_distribute_reduction_codegen.cpp
View File

@ -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