This is an archive of the discontinued LLVM Phabricator instance.

Differential D99622

[OpenMP51] Accept `primary` as proc bind affinity policy in Clang
ClosedPublic

Authored by cchen on Mar 30 2021, 3:03 PM.

Details

Reviewers
ABataev
jdoerfert
Commits
rGcba422264c7f: [OpenMP51] Accept `primary` as proc bind affinity policy in Clang

Diff Detail

Unit TestsFailed

TimeTest
310 msx64 debian > Clang.OpenMP::teams_distribute_parallel_for_proc_bind_codegen.cpp
280 msx64 debian > Clang.OpenMP::teams_distribute_parallel_for_simd_proc_bind_codegen.cpp
600 msx64 windows > Clang.OpenMP::teams_distribute_parallel_for_proc_bind_codegen.cpp
720 msx64 windows > Clang.OpenMP::teams_distribute_parallel_for_simd_proc_bind_codegen.cpp

Event Timeline

cchen created this revision.Mar 30 2021, 3:03 PM
Herald added subscribers: guansong, yaxunl. · View Herald TranscriptMar 30 2021, 3:03 PM
cchen requested review of this revision.Mar 30 2021, 3:03 PM
Herald added a reviewer: jdoerfert. · View Herald TranscriptMar 30 2021, 3:03 PM
Herald added projects: Restricted Project, Restricted Project. · View Herald Transcript
Herald added subscribers: llvm-commits, cfe-commits, sstefan1. · View Herald Transcript
Harbormaster completed remote builds in B96407: Diff 334279.Mar 30 2021, 3:46 PM
ABataev added a comment.Mar 31 2021, 4:31 AM

Could you add a small codegen test?

cchen updated this revision to Diff 334482.Mar 31 2021, 10:36 AM

Add some codegen tests

ABataev added inline comments.Mar 31 2021, 10:41 AM
clang/test/OpenMP/distribute_parallel_for_proc_bind_codegen.cpp
108–135

Same, could you move it to a different file?

clang/test/OpenMP/nvptx_target_parallel_proc_bind_codegen.cpp
110–147

Put this part in a separate file, please

cchen added inline comments.Mar 31 2021, 10:56 AM
clang/test/OpenMP/nvptx_target_parallel_proc_bind_codegen.cpp
110–147

Do you mean I should put all the primary tests into a single file and don't do any changes to the existing proc_bind tests?

ABataev added inline comments.Mar 31 2021, 11:00 AM
clang/test/OpenMP/nvptx_target_parallel_proc_bind_codegen.cpp
110–147

Yes, keep the original tests and add new ones.

Harbormaster completed remote builds in B96560: Diff 334482.Mar 31 2021, 11:31 AM
cchen updated this revision to Diff 334720.Apr 1 2021, 8:55 AM

Update codegen test

ABataev accepted this revision.Apr 1 2021, 8:59 AM

LG

This revision is now accepted and ready to land.Apr 1 2021, 8:59 AM
Harbormaster completed remote builds in B96742: Diff 334720.Apr 1 2021, 9:52 AM
cchen updated this revision to Diff 334844.Apr 1 2021, 4:06 PM

Rebase

This revision was landed with ongoing or failed builds.Apr 1 2021, 4:13 PM
Closed by commit rGcba422264c7f: [OpenMP51] Accept `primary` as proc bind affinity policy in Clang (authored by cchen). · Explain Why
This revision was automatically updated to reflect the committed changes.
cchen added a commit: rGcba422264c7f: [OpenMP51] Accept `primary` as proc bind affinity policy in Clang.
Harbormaster completed remote builds in B96824: Diff 334844.Apr 1 2021, 5:10 PM

Revision Contents

Diff 334482

clang/lib/Sema/SemaOpenMP.cpp

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 14,063 Lines • ▼ Show 20 LinesOMPClause *Sema::ActOnOpenMPProcBindClause(ProcBindKind Kind,
SourceLocation KindKwLoc, SourceLocation KindKwLoc,
SourceLocation StartLoc, SourceLocation StartLoc,
SourceLocation LParenLoc, SourceLocation LParenLoc,
SourceLocation EndLoc) { SourceLocation EndLoc) {
if (Kind == OMP_PROC_BIND_unknown) { if (Kind == OMP_PROC_BIND_unknown) {
Diag(KindKwLoc, diag::err_omp_unexpected_clause_value) Diag(KindKwLoc, diag::err_omp_unexpected_clause_value)
<< getListOfPossibleValues(OMPC_proc_bind, << getListOfPossibleValues(OMPC_proc_bind,
/*First=*/unsigned(OMP_PROC_BIND_master), /*First=*/unsigned(OMP_PROC_BIND_master),
/*Last=*/5) /*Last=*/
unsigned(LangOpts.OpenMP > 50
? OMP_PROC_BIND_primary
: OMP_PROC_BIND_spread) +
1)
<< getOpenMPClauseName(OMPC_proc_bind); << getOpenMPClauseName(OMPC_proc_bind);
return nullptr; return nullptr;
} }
if (Kind == OMP_PROC_BIND_primary && LangOpts.OpenMP < 51)
Diag(KindKwLoc, diag::err_omp_unexpected_clause_value)
<< getListOfPossibleValues(OMPC_proc_bind,
/*First=*/unsigned(OMP_PROC_BIND_master),
/*Last=*/
unsigned(OMP_PROC_BIND_spread) + 1)
<< getOpenMPClauseName(OMPC_proc_bind);
return new (Context) return new (Context)
OMPProcBindClause(Kind, KindKwLoc, StartLoc, LParenLoc, EndLoc); OMPProcBindClause(Kind, KindKwLoc, StartLoc, LParenLoc, EndLoc);
} }
OMPClause *Sema::ActOnOpenMPAtomicDefaultMemOrderClause( OMPClause *Sema::ActOnOpenMPAtomicDefaultMemOrderClause(
OpenMPAtomicDefaultMemOrderClauseKind Kind, SourceLocation KindKwLoc, OpenMPAtomicDefaultMemOrderClauseKind Kind, SourceLocation KindKwLoc,
SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc) { SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc) {
if (Kind == OMPC_ATOMIC_DEFAULT_MEM_ORDER_unknown) { if (Kind == OMPC_ATOMIC_DEFAULT_MEM_ORDER_unknown) {
▲ Show 20 Lines • Show All 6,513 Lines • Show Last 20 Lines

clang/test/OpenMP/distribute_parallel_for_proc_bind_codegen.cpp

// add -fopenmp-targets // add -fopenmp-targets
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
// expected-no-diagnostics // expected-no-diagnostics
#ifndef HEADER #ifndef HEADER
#define HEADER #define HEADER
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK1
typedef __INTPTR_TYPE__ intptr_t; typedef __INTPTR_TYPE__ intptr_t;
// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* } // CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" // CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } // CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo(); void foo();
struct S { struct S {
intptr_t a, b, c; intptr_t a, b, c;
S(intptr_t a) : a(a) {} S(intptr_t a) : a(a) {}
operator char() { return a; } operator char() { return a; }
~S() {} ~S() {}
}; };
template <typename T> template <typename T>
T tmain() { T tmain() {
#pragma omp target #pragma omp target
#pragma omp teams #pragma omp teams
#pragma omp distribute parallel for proc_bind(master) #pragma omp distribute parallel for proc_bind(master)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
return T(); return T();
} }
int main() { int main() {
// CHECK-LABEL: @main // CK1-LABEL: @main
#pragma omp target
#pragma omp teams
#pragma omp distribute parallel for proc_bind(spread)
for(int i = 0; i < 1000; i++) {}
#pragma omp target
#pragma omp teams
#pragma omp distribute parallel for proc_bind(close)
for(int i = 0; i < 1000; i++) {}
return tmain<int>();
}
// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK1: call void [[OFFL1:@.+]]()
// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK1: call void [[OFFL2:@.+]]()
// CK1: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
// CK1: ret i32 [[CALL_RET]]
// CK1: define{{.+}} void [[OFFL1]](
// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK1: ret void
// CK1: define{{.+}} [[OFFL2]]()
// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK1: ret void
// CK1: define{{.+}} [[TMAIN]]()
// CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK1: call void [[OFFL3:@.+]]()
// CK1: define{{.+}} [[OFFL3]]()
// CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
// CK1: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
// CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
// CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK1: ret void
#endif
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK2
typedef __INTPTR_TYPE__ intptr_t;
// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo();
struct S {
intptr_t a, b, c;
S(intptr_t a) : a(a) {}
operator char() { return a; }
~S() {}
};
template <typename T>
T tmain() {
#pragma omp target
#pragma omp teams
#pragma omp distribute parallel for proc_bind(primary)
for(int i = 0; i < 1000; i++) {}
return T();
}
int main() {
// CK2-LABEL: @main
ABataevUnsubmitted
Not Done
ReplyInline Actions

Same, could you move it to a different file?

ABataev: Same, could you move it to a different file?
#pragma omp target #pragma omp target
#pragma omp teams #pragma omp teams
#pragma omp distribute parallel for proc_bind(spread) #pragma omp distribute parallel for proc_bind(spread)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
#pragma omp target #pragma omp target
#pragma omp teams #pragma omp teams
#pragma omp distribute parallel for proc_bind(close) #pragma omp distribute parallel for proc_bind(close)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
return tmain<int>(); return tmain<int>();
} }
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL1:@.+]]() // CK2: call void [[OFFL1:@.+]]()
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL2:@.+]]() // CK2: call void [[OFFL2:@.+]]()
// CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]() // CK2: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
// CHECK: ret i32 [[CALL_RET]] // CK2: ret i32 [[CALL_RET]]
// CHECK: define{{.+}} void [[OFFL1]]( // CK2: define{{.+}} void [[OFFL1]](
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}}) // CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]], // CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4) // CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK2: ret void
// CHECK: define{{.+}} [[OFFL2]]() // CK2: define{{.+}} [[OFFL2]]()
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}}) // CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]], // CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3) // CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK2: ret void
// CHECK: define{{.+}} [[TMAIN]]() // CK2: define{{.+}} [[TMAIN]]()
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL3:@.+]]() // CK2: call void [[OFFL3:@.+]]()
// CHECK: define{{.+}} [[OFFL3]]() // CK2: define{{.+}} [[OFFL3]]()
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}}) // CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]], // CK2: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2) // CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK2: ret void
#endif
#endif #endif

clang/test/OpenMP/nvptx_target_parallel_proc_bind_codegen.cpp

// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
// expected-no-diagnostics // expected-no-diagnostics
#ifndef HEADER #ifndef HEADER
#define HEADER #define HEADER
// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32
#ifdef CK1
// Check that the execution mode of all 3 target regions on the gpu is set to SPMD Mode. // Check that the execution mode of all 3 target regions on the gpu is set to SPMD Mode.
// CHECK-DAG: {{@__omp_offloading_.+l29}}_exec_mode = weak constant i8 0 // CK1-DAG: {{@__omp_offloading_.+l29}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l33}}_exec_mode = weak constant i8 0 // CK1-DAG: {{@__omp_offloading_.+l33}}_exec_mode = weak constant i8 0
// CHECK-DAG: {{@__omp_offloading_.+l38}}_exec_mode = weak constant i8 0 // CK1-DAG: {{@__omp_offloading_.+l38}}_exec_mode = weak constant i8 0
template<typename tx> template<typename tx>
tx ftemplate(int n) { tx ftemplate(int n) {
tx a = 0; tx a = 0;
short aa = 0; short aa = 0;
tx b[10]; tx b[10];
#pragma omp target parallel proc_bind(master) #pragma omp target parallel proc_bind(master)
Show All 18 Lines
int bar(int n){ int bar(int n){
int a = 0; int a = 0;
a += ftemplate<int>(n); a += ftemplate<int>(n);
return a; return a;
} }
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l29}}( // CK1-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l29}}(
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1) // CK1: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
// CHECK: call void @__kmpc_data_sharing_init_stack_spmd // CK1: call void @__kmpc_data_sharing_init_stack_spmd
// CHECK: br label {{%?}}[[EXEC:.+]] // CK1: br label {{%?}}[[EXEC:.+]]
// //
// CHECK: [[EXEC]] // CK1: [[EXEC]]
// CHECK-NOT: call void @__kmpc_push_proc_bind // CK1-NOT: call void @__kmpc_push_proc_bind
// CHECK: {{call|invoke}} void [[OP1:@.+]]( // CK1: {{call|invoke}} void [[OP1:@.+]](
// CHECK: br label {{%?}}[[DONE:.+]] // CK1: br label {{%?}}[[DONE:.+]]
// //
// CHECK: [[DONE]] // CK1: [[DONE]]
// CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 1) // CK1: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
// CHECK: br label {{%?}}[[EXIT:.+]] // CK1: br label {{%?}}[[EXIT:.+]]
// //
// CHECK: [[EXIT]] // CK1: [[EXIT]]
// CHECK: ret void // CK1: ret void
// CHECK: } // CK1: }
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l33}}( // CK1-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l33}}(
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1) // CK1: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
// CHECK: call void @__kmpc_data_sharing_init_stack_spmd // CK1: call void @__kmpc_data_sharing_init_stack_spmd
// CHECK: br label {{%?}}[[EXEC:.+]] // CK1: br label {{%?}}[[EXEC:.+]]
// //
// CHECK: [[EXEC]] // CK1: [[EXEC]]
// CHECK-NOT: call void @__kmpc_push_proc_bind // CK1-NOT: call void @__kmpc_push_proc_bind
// CHECK: {{call|invoke}} void [[OP1:@.+]]( // CK1: {{call|invoke}} void [[OP1:@.+]](
// CHECK: br label {{%?}}[[DONE:.+]] // CK1: br label {{%?}}[[DONE:.+]]
// //
// CHECK: [[DONE]] // CK1: [[DONE]]
// CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 1) // CK1: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
// CHECK: br label {{%?}}[[EXIT:.+]] // CK1: br label {{%?}}[[EXIT:.+]]
// //
// CHECK: [[EXIT]] // CK1: [[EXIT]]
// CHECK: ret void // CK1: ret void
// CHECK: } // CK1: }
// CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l38}}( // CK1-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l38}}(
// CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1) // CK1: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
// CHECK: call void @__kmpc_data_sharing_init_stack_spmd // CK1: call void @__kmpc_data_sharing_init_stack_spmd
// CHECK: br label {{%?}}[[EXEC:.+]] // CK1: br label {{%?}}[[EXEC:.+]]
// //
// CHECK: [[EXEC]] // CK1: [[EXEC]]
// CHECK-NOT: call void @__kmpc_push_proc_bind // CK1-NOT: call void @__kmpc_push_proc_bind
// CHECK: {{call|invoke}} void [[OP1:@.+]]( // CK1: {{call|invoke}} void [[OP1:@.+]](
// CHECK: br label {{%?}}[[DONE:.+]] // CK1: br label {{%?}}[[DONE:.+]]
// //
// CHECK: [[DONE]] // CK1: [[DONE]]
// CHECK: call void @__kmpc_spmd_kernel_deinit_v2(i16 1) // CK1: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
// CHECK: br label {{%?}}[[EXIT:.+]] // CK1: br label {{%?}}[[EXIT:.+]]
// //
// CHECK: [[EXIT]] // CK1: [[EXIT]]
// CHECK: ret void // CK1: ret void
// CHECK: } // CK1: }
#endif
// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=45 -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-cuda-mode -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-cuda-mode -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32
#ifdef CK2
// Check that the execution mode of all 3 target regions on the gpu is set to SPMD Mode.
// CK2-DAG: {{@__omp_offloading_.+l132}}_exec_mode = weak constant i8 0
template<typename tx>
tx ftemplate(int n) {
tx a = 0;
short aa = 0;
tx b[10];
#pragma omp target parallel proc_bind(master)
{
}
return a;
}
int bar(int n){
int a = 0;
a += ftemplate<int>(n);
return a;
}
// CK2-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l132}}(
ABataevUnsubmitted
Not Done
ReplyInline Actions

Put this part in a separate file, please

ABataev: Put this part in a separate file, please
cchenAuthorUnsubmitted
Done
ReplyInline Actions

Do you mean I should put all the primary tests into a single file and don't do any changes to the existing proc_bind tests?

cchen: Do you mean I should put all the `primary` tests into a single file and don't do any changes to…
ABataevUnsubmitted
Not Done
ReplyInline Actions

Yes, keep the original tests and add new ones.

ABataev: Yes, keep the original tests and add new ones.
// CK2: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1)
// CK2: call void @__kmpc_data_sharing_init_stack_spmd
// CK2: br label {{%?}}[[EXEC:.+]]
//
// CK2: [[EXEC]]
// CK2-NOT: call void @__kmpc_push_proc_bind
// CK2: {{call|invoke}} void [[OP1:@.+]](
// CK2: br label {{%?}}[[DONE:.+]]
//
// CK2: [[DONE]]
// CK2: call void @__kmpc_spmd_kernel_deinit_v2(i16 1)
// CK2: br label {{%?}}[[EXIT:.+]]
//
// CK2: [[EXIT]]
// CK2: ret void
// CK2: }
#endif
#endif #endif

clang/test/OpenMP/parallel_ast_print.cpp

// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s // RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s // RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s // RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s // RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
// expected-no-diagnostics // expected-no-diagnostics
#ifndef HEADER #ifndef HEADER
#define HEADER #define HEADER
void foo() {} void foo() {}
struct S1 { struct S1 {
▲ Show 20 Lines • Show All 131 Lines • ▼ Show 20 LinesT tmain(T argc, T *argv) {
T b = argc, c, d, e, f, g; T b = argc, c, d, e, f, g;
static T a; static T a;
S<T> s; S<T> s;
T arr[C][10], arr1[C]; T arr[C][10], arr1[C];
#pragma omp parallel #pragma omp parallel
a=2; a=2;
#pragma omp parallel default(none), private(argc,b) firstprivate(argv) shared (d) if (parallel:argc > 0) num_threads(C) copyin(S<T>::TS, thrp) proc_bind(master) reduction(+:c, arr1[argc]) reduction(max:e, arr[:C][0:10]) #pragma omp parallel default(none), private(argc,b) firstprivate(argv) shared (d) if (parallel:argc > 0) num_threads(C) copyin(S<T>::TS, thrp) proc_bind(master) reduction(+:c, arr1[argc]) reduction(max:e, arr[:C][0:10])
foo(); foo();
#ifdef OMP51
#pragma omp parallel default(none), private(argc,b) firstprivate(argv) shared (d) if (parallel:argc > 0) num_threads(C) copyin(S<T>::TS, thrp) proc_bind(primary) reduction(+:c, arr1[argc]) reduction(max:e, arr[:C][0:10])
foo();
#endif
#pragma omp parallel if (C) num_threads(s) proc_bind(close) reduction(^:e, f, arr[0:C][:argc]) reduction(default, && : g) reduction(task,+:argc) #pragma omp parallel if (C) num_threads(s) proc_bind(close) reduction(^:e, f, arr[0:C][:argc]) reduction(default, && : g) reduction(task,+:argc)
foo(); foo();
return 0; return 0;
} }
// CHECK: template <typename T, int C> T tmain(T argc, T *argv) { // CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
// CHECK-NEXT: T b = argc, c, d, e, f, g; // CHECK-NEXT: T b = argc, c, d, e, f, g;
// CHECK-NEXT: static T a; // CHECK-NEXT: static T a;
// CHECK-NEXT: S<T> s; // CHECK-NEXT: S<T> s;
// CHECK-NEXT: T arr[C][10], arr1[C]; // CHECK-NEXT: T arr[C][10], arr1[C];
// CHECK-NEXT: #pragma omp parallel{{$}} // CHECK-NEXT: #pragma omp parallel{{$}}
// CHECK-NEXT: a = 2; // CHECK-NEXT: a = 2;
// CHECK-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(C) copyin(S<T>::TS,thrp) proc_bind(master) reduction(+: c,arr1[argc]) reduction(max: e,arr[:C][0:10]) // CHECK-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(C) copyin(S<T>::TS,thrp) proc_bind(master) reduction(+: c,arr1[argc]) reduction(max: e,arr[:C][0:10])
// CHECK-NEXT: foo() // CHECK-NEXT: foo()
// OMP51-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(C) copyin(S<T>::TS,thrp) proc_bind(primary) reduction(+: c,arr1[argc]) reduction(max: e,arr[:C][0:10])
// OMP51-NEXT: foo()
// CHECK-NEXT: #pragma omp parallel if(C) num_threads(s) proc_bind(close) reduction(^: e,f,arr[0:C][:argc]) reduction(default, &&: g) reduction(task, +: argc) // CHECK-NEXT: #pragma omp parallel if(C) num_threads(s) proc_bind(close) reduction(^: e,f,arr[0:C][:argc]) reduction(default, &&: g) reduction(task, +: argc)
// CHECK-NEXT: foo() // CHECK-NEXT: foo()
// CHECK: template<> int tmain<int, 5>(int argc, int *argv) { // CHECK: template<> int tmain<int, 5>(int argc, int *argv) {
// CHECK-NEXT: int b = argc, c, d, e, f, g; // CHECK-NEXT: int b = argc, c, d, e, f, g;
// CHECK-NEXT: static int a; // CHECK-NEXT: static int a;
// CHECK-NEXT: S<int> s; // CHECK-NEXT: S<int> s;
// CHECK-NEXT: int arr[5][10], arr1[5]; // CHECK-NEXT: int arr[5][10], arr1[5];
// CHECK-NEXT: #pragma omp parallel // CHECK-NEXT: #pragma omp parallel
// CHECK-NEXT: a = 2; // CHECK-NEXT: a = 2;
// CHECK-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(5) copyin(S<int>::TS,thrp) proc_bind(master) reduction(+: c,arr1[argc]) reduction(max: e,arr[:5][0:10]) // CHECK-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(5) copyin(S<int>::TS,thrp) proc_bind(master) reduction(+: c,arr1[argc]) reduction(max: e,arr[:5][0:10])
// CHECK-NEXT: foo() // CHECK-NEXT: foo()
// OMP51-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(5) copyin(S<int>::TS,thrp) proc_bind(primary) reduction(+: c,arr1[argc]) reduction(max: e,arr[:5][0:10])
// OMP51-NEXT: foo()
// CHECK-NEXT: #pragma omp parallel if(5) num_threads(s) proc_bind(close) reduction(^: e,f,arr[0:5][:argc]) reduction(default, &&: g) reduction(task, +: argc) // CHECK-NEXT: #pragma omp parallel if(5) num_threads(s) proc_bind(close) reduction(^: e,f,arr[0:5][:argc]) reduction(default, &&: g) reduction(task, +: argc)
// CHECK-NEXT: foo() // CHECK-NEXT: foo()
// CHECK: template<> long tmain<long, 1>(long argc, long *argv) { // CHECK: template<> long tmain<long, 1>(long argc, long *argv) {
// CHECK-NEXT: long b = argc, c, d, e, f, g; // CHECK-NEXT: long b = argc, c, d, e, f, g;
// CHECK-NEXT: static long a; // CHECK-NEXT: static long a;
// CHECK-NEXT: S<long> s; // CHECK-NEXT: S<long> s;
// CHECK-NEXT: long arr[1][10], arr1[1]; // CHECK-NEXT: long arr[1][10], arr1[1];
// CHECK-NEXT: #pragma omp parallel // CHECK-NEXT: #pragma omp parallel
// CHECK-NEXT: a = 2; // CHECK-NEXT: a = 2;
// CHECK-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(1) copyin(S<long>::TS,thrp) proc_bind(master) reduction(+: c,arr1[argc]) reduction(max: e,arr[:1][0:10]) // CHECK-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(1) copyin(S<long>::TS,thrp) proc_bind(master) reduction(+: c,arr1[argc]) reduction(max: e,arr[:1][0:10])
// CHECK-NEXT: foo() // CHECK-NEXT: foo()
// OMP51-NEXT: #pragma omp parallel default(none) private(argc,b) firstprivate(argv) shared(d) if(parallel: argc > 0) num_threads(1) copyin(S<long>::TS,thrp) proc_bind(primary) reduction(+: c,arr1[argc]) reduction(max: e,arr[:1][0:10])
// OMP51-NEXT: foo()
// CHECK-NEXT: #pragma omp parallel if(1) num_threads(s) proc_bind(close) reduction(^: e,f,arr[0:1][:argc]) reduction(default, &&: g) reduction(task, +: argc) // CHECK-NEXT: #pragma omp parallel if(1) num_threads(s) proc_bind(close) reduction(^: e,f,arr[0:1][:argc]) reduction(default, &&: g) reduction(task, +: argc)
// CHECK-NEXT: foo() // CHECK-NEXT: foo()
enum Enum { }; enum Enum { };
int main (int argc, char **argv) { int main (int argc, char **argv) {
long x; long x;
int b = argc, c, d, e, f, g; int b = argc, c, d, e, f, g;
▲ Show 20 Lines • Show All 42 Lines • Show Last 20 Lines

clang/test/OpenMP/parallel_master_codegen.cpp

Show First 20 Lines • Show All 642 Lines • ▼ Show 20 Lines
// CK9-NOT: __kmpc_global_thread_num // CK9-NOT: __kmpc_global_thread_num
// CK9: call i32 @__kmpc_master({{.+}}) // CK9: call i32 @__kmpc_master({{.+}})
// CK9: [[FOUR:%.+]] = load i32, i32* [[A_FP_ADDR]] // CK9: [[FOUR:%.+]] = load i32, i32* [[A_FP_ADDR]]
// CK9: [[INC:%.+]] = add nsw i32 [[FOUR]] // CK9: [[INC:%.+]] = add nsw i32 [[FOUR]]
// CK9: store i32 [[INC]], i32* [[A_FP_ADDR]] // CK9: store i32 [[INC]], i32* [[A_FP_ADDR]]
// CK9-NOT: __kmpc_global_thread_num // CK9-NOT: __kmpc_global_thread_num
// CK9: call void @__kmpc_end_master({{.+}}) // CK9: call void @__kmpc_end_master({{.+}})
#endif #endif
#ifdef CK10
///==========================================================================///
// RUN: %clang_cc1 -DCK10 -verify -fopenmp -fopenmp-version=51 -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s --check-prefix CK10
// RUN: %clang_cc1 -DCK10 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK10 -fopenmp -fopenmp-version=51 -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK10
// RUN: %clang_cc1 -DCK10 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -triple x86_64-unknown-unknown -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -DCK10 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK10 -fopenmp-simd -fopenmp-version=51 -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
typedef __INTPTR_TYPE__ intptr_t;
// CK10-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CK10-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CK10-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo();
struct S {
intptr_t a, b, c;
S(intptr_t a) : a(a) {}
operator char() { return a; }
~S() {}
};
template <typename T>
T tmain() {
#pragma omp parallel master proc_bind(primary)
foo();
return T();
}
int main() {
#pragma omp parallel master proc_bind(spread)
foo();
#pragma omp parallel master proc_bind(close)
foo();
return tmain<int>();
}
// CK10-LABEL: @main
// CK10: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
// CK10: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 4)
// CK10: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK10: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 3)
// CK10: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK10-LABEL: @{{.+}}tmain
// CK10: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
// CK10: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 5)
// CK10: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK10: ret i32 0
// CK10-NEXT: }
// CK10: call i32 @__kmpc_master(%struct.ident_t* [[DEF_LOC_2]], i32 [[ONE:%.+]])
// CK10: call void @__kmpc_end_master(%struct.ident_t* [[DEF_LOC_2]], i32 [[ONE]])
#endif
#endif #endif

clang/test/OpenMP/parallel_proc_bind_codegen.cpp

// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
// expected-no-diagnostics // expected-no-diagnostics
#ifndef HEADER #ifndef HEADER
#define HEADER #define HEADER
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK1
typedef __INTPTR_TYPE__ intptr_t; typedef __INTPTR_TYPE__ intptr_t;
// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* } // CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" // CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } // CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo(); void foo();
struct S { struct S {
intptr_t a, b, c; intptr_t a, b, c;
S(intptr_t a) : a(a) {} S(intptr_t a) : a(a) {}
operator char() { return a; } operator char() { return a; }
~S() {} ~S() {}
Show All 9 Lines
int main() { int main() {
#pragma omp parallel proc_bind(spread) #pragma omp parallel proc_bind(spread)
foo(); foo();
#pragma omp parallel proc_bind(close) #pragma omp parallel proc_bind(close)
foo(); foo();
return tmain<int>(); return tmain<int>();
} }
// CHECK-LABEL: @main // CK1-LABEL: @main
// CHECK: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]]) // CK1: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 4) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 4)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 3) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 3)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK-LABEL: @{{.+}}tmain // CK1-LABEL: @{{.+}}tmain
// CHECK: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]]) // CK1: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 2) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 2)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret i32 0 // CK1: ret i32 0
// CHECK-NEXT: } // CK1-NEXT: }
#endif
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -triple %itanium_abi_triple -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK2
typedef __INTPTR_TYPE__ intptr_t;
// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo();
struct S {
intptr_t a, b, c;
S(intptr_t a) : a(a) {}
operator char() { return a; }
~S() {}
};
template <typename T>
T tmain() {
#pragma omp parallel proc_bind(primary)
foo();
return T();
}
int main() {
#pragma omp parallel proc_bind(primary)
foo();
return tmain<int>();
}
// CK2-LABEL: @main
// CK2: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 5)
// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK2-LABEL: @{{.+}}tmain
// CK2: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEF_LOC_2]])
// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID]], i32 5)
// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK2: ret i32 0
// CK2-NEXT: }
#endif
#endif #endif

clang/test/OpenMP/parallel_proc_bind_messages.cpp

// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 -o - %s -Wuninitialized // RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 -o - %s -Wuninitialized // RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
void foo(); void foo();
int main(int argc, char **argv) { int main(int argc, char **argv) {
#pragma omp parallel proc_bind // expected-error {{expected '(' after 'proc_bind'}} #pragma omp parallel proc_bind // expected-error {{expected '(' after 'proc_bind'}}
#pragma omp parallel proc_bind ( // expected-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}} expected-error {{expected ')'}} expected-note {{to match this '('}} #pragma omp parallel proc_bind ( // ge51-error {{expected 'master', 'close', 'spread' or 'primary' in OpenMP clause 'proc_bind'}} lt51-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
#pragma omp parallel proc_bind () // expected-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}} #pragma omp parallel proc_bind () // ge51-error {{expected 'master', 'close', 'spread' or 'primary' in OpenMP clause 'proc_bind'}} lt51-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}}
#pragma omp parallel proc_bind (master // expected-error {{expected ')'}} expected-note {{to match this '('}} #pragma omp parallel proc_bind (master // expected-error {{expected ')'}} expected-note {{to match this '('}}
#pragma omp parallel proc_bind (close), proc_bind(spread) // expected-error {{directive '#pragma omp parallel' cannot contain more than one 'proc_bind' clause}} #pragma omp parallel proc_bind (close), proc_bind(spread) // expected-error {{directive '#pragma omp parallel' cannot contain more than one 'proc_bind' clause}}
#pragma omp parallel proc_bind (x) // expected-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}} #pragma omp parallel proc_bind (x) // ge51-error {{expected 'master', 'close', 'spread' or 'primary' in OpenMP clause 'proc_bind'}} lt51-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}}
foo(); foo();
#pragma omp parallel proc_bind(master) #pragma omp parallel proc_bind(master)
++argc; ++argc;
#pragma omp parallel proc_bind(close) #pragma omp parallel proc_bind(close)
#pragma omp parallel proc_bind(spread) #pragma omp parallel proc_bind(spread)
++argc; ++argc;
#pragma omp parallel proc_bind(primary) // lt51-error {{expected 'master', 'close' or 'spread' in OpenMP clause 'proc_bind'}}
++argc;
return 0; return 0;
} }

clang/test/OpenMP/target_teams_distribute_parallel_for_proc_bind_codegen.cpp

// add -fopenmp-targets // add -fopenmp-targets
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
// expected-no-diagnostics // expected-no-diagnostics
#ifndef HEADER #ifndef HEADER
#define HEADER #define HEADER
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
#ifdef CK1
typedef __INTPTR_TYPE__ intptr_t; typedef __INTPTR_TYPE__ intptr_t;
// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* } // CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" // CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } // CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo(); void foo();
struct S { struct S {
intptr_t a, b, c; intptr_t a, b, c;
S(intptr_t a) : a(a) {} S(intptr_t a) : a(a) {}
operator char() { return a; } operator char() { return a; }
~S() {} ~S() {}
}; };
template <typename T> template <typename T>
T tmain() { T tmain() {
#pragma omp target teams distribute parallel for proc_bind(master) #pragma omp target teams distribute parallel for proc_bind(master)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
return T(); return T();
} }
int main() { int main() {
// CHECK-LABEL: @main // CK1-LABEL: @main
#pragma omp target teams distribute parallel for proc_bind(spread) #pragma omp target teams distribute parallel for proc_bind(spread)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
#pragma omp target teams distribute parallel for proc_bind(close) #pragma omp target teams distribute parallel for proc_bind(close)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
return tmain<int>(); return tmain<int>();
} }
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL1:@.+]]() // CK1: call void [[OFFL1:@.+]]()
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL2:@.+]]() // CK1: call void [[OFFL2:@.+]]()
// CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]() // CK1: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
// CHECK: ret i32 [[CALL_RET]] // CK1: ret i32 [[CALL_RET]]
// CHECK: define{{.+}} void [[OFFL1]]( // CK1: define{{.+}} void [[OFFL1]](
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
// CHECK: define{{.+}} [[OFFL2]]() // CK1: define{{.+}} [[OFFL2]]()
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
// CHECK: define{{.+}} [[TMAIN]]() // CK1: define{{.+}} [[TMAIN]]()
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL3:@.+]]() // CK1: call void [[OFFL3:@.+]]()
// CHECK: define{{.+}} [[OFFL3]]() // CK1: define{{.+}} [[OFFL3]]()
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
#endif
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
#ifdef CK2
typedef __INTPTR_TYPE__ intptr_t;
// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo();
template <typename T>
T tmain() {
#pragma omp target teams distribute parallel for proc_bind(primary)
for(int i = 0; i < 1000; i++) {}
return T();
}
int main() {
// CK2-LABEL: @main
#pragma omp target teams distribute parallel for proc_bind(primary)
for(int i = 0; i < 1000; i++) {}
return tmain<int>();
}
// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK2: call void [[OFFL1:@.+]]()
// CK2: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
// CK2: ret i32 [[CALL_RET]]
// CK2: define{{.+}} void [[OFFL1]](
// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK2: ret void
// CK2: define{{.+}} [[TMAIN]]()
// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK2: call void [[OFFL2:@.+]]()
// CK2: define{{.+}} [[OFFL2]]()
// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
// CK2: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK2: ret void
#endif
#endif #endif

clang/test/OpenMP/target_teams_distribute_parallel_for_simd_proc_bind_codegen.cpp

// add -fopenmp-targets // add -fopenmp-targets
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
// expected-no-diagnostics // expected-no-diagnostics
#ifndef HEADER #ifndef HEADER
#define HEADER #define HEADER
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
#ifdef CK1
typedef __INTPTR_TYPE__ intptr_t; typedef __INTPTR_TYPE__ intptr_t;
// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* } // CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" // CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } // CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo(); void foo();
struct S { struct S {
intptr_t a, b, c; intptr_t a, b, c;
S(intptr_t a) : a(a) {} S(intptr_t a) : a(a) {}
operator char() { return a; } operator char() { return a; }
~S() {} ~S() {}
}; };
template <typename T> template <typename T>
T tmain() { T tmain() {
#pragma omp target teams distribute parallel for simd proc_bind(master) #pragma omp target teams distribute parallel for simd proc_bind(master)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
return T(); return T();
} }
int main() { int main() {
// CHECK-LABEL: @main // CK1-LABEL: @main
#pragma omp target teams distribute parallel for simd proc_bind(spread) #pragma omp target teams distribute parallel for simd proc_bind(spread)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
#pragma omp target teams distribute parallel for simd proc_bind(close) #pragma omp target teams distribute parallel for simd proc_bind(close)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
return tmain<int>(); return tmain<int>();
} }
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL1:@.+]]() // CK1: call void [[OFFL1:@.+]]()
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL2:@.+]]() // CK1: call void [[OFFL2:@.+]]()
// CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]() // CK1: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
// CHECK: ret i32 [[CALL_RET]] // CK1: ret i32 [[CALL_RET]]
// CHECK: define{{.+}} void [[OFFL1]]( // CK1: define{{.+}} void [[OFFL1]](
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
// CHECK: define{{.+}} [[OFFL2]]() // CK1: define{{.+}} [[OFFL2]]()
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
// CHECK: define{{.+}} [[TMAIN]]() // CK1: define{{.+}} [[TMAIN]]()
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL3:@.+]]() // CK1: call void [[OFFL3:@.+]]()
// CHECK: define{{.+}} [[OFFL3]]() // CK1: define{{.+}} [[OFFL3]]()
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
#endif
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY
// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
#ifdef CK2
typedef __INTPTR_TYPE__ intptr_t;
// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo();
template <typename T>
T tmain() {
#pragma omp target teams distribute parallel for simd proc_bind(primary)
for(int i = 0; i < 1000; i++) {}
return T();
}
int main() {
// CK2-LABEL: @main
#pragma omp target teams distribute parallel for simd proc_bind(primary)
for(int i = 0; i < 1000; i++) {}
return tmain<int>();
}
// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK2: call void [[OFFL1:@.+]]()
// CK2: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
// CK2: ret i32 [[CALL_RET]]
// CK2: define{{.+}} void [[OFFL1]](
// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK2: ret void
// CK2: define{{.+}} [[TMAIN]]()
// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK2: call void [[OFFL2:@.+]]()
// CK2: define{{.+}} [[OFFL2]]()
// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
// CK2: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK2: ret void
#endif
#endif #endif

clang/test/OpenMP/teams_distribute_parallel_for_proc_bind_codegen.cpp

// add -fopenmp-targets // add -fopenmp-targets
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
// expected-no-diagnostics // expected-no-diagnostics
#ifndef HEADER #ifndef HEADER
#define HEADER #define HEADER
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK1
typedef __INTPTR_TYPE__ intptr_t; typedef __INTPTR_TYPE__ intptr_t;
// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* } // CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" // CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } // CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo(); void foo();
struct S { struct S {
intptr_t a, b, c; intptr_t a, b, c;
S(intptr_t a) : a(a) {} S(intptr_t a) : a(a) {}
operator char() { return a; } operator char() { return a; }
~S() {} ~S() {}
}; };
template <typename T> template <typename T>
T tmain() { T tmain() {
#pragma omp target #pragma omp target
#pragma omp teams distribute parallel for proc_bind(master) #pragma omp teams distribute parallel for proc_bind(master)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
return T(); return T();
} }
int main() { int main() {
// CHECK-LABEL: @main // CK1-LABEL: @main
#pragma omp target #pragma omp target
#pragma omp teams distribute parallel for proc_bind(spread) #pragma omp teams distribute parallel for proc_bind(spread)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
#pragma omp target #pragma omp target
#pragma omp teams distribute parallel for proc_bind(close) #pragma omp teams distribute parallel for proc_bind(close)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
return tmain<int>(); return tmain<int>();
} }
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL1:@.+]]() // CK1: call void [[OFFL1:@.+]]()
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL2:@.+]]() // CK1: call void [[OFFL2:@.+]]()
// CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]() // CK1: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
// CHECK: ret i32 [[CALL_RET]] // CK1: ret i32 [[CALL_RET]]
// CHECK: define{{.+}} void [[OFFL1]]( // CK1: define{{.+}} void [[OFFL1]](
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
// CHECK: define{{.+}} [[OFFL2]]() // CK1: define{{.+}} [[OFFL2]]()
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
// CHECK: define{{.+}} [[TMAIN]]() // CK1: define{{.+}} [[TMAIN]]()
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL3:@.+]]() // CK1: call void [[OFFL3:@.+]]()
// CHECK: define{{.+}} [[OFFL3]]() // CK1: define{{.+}} [[OFFL3]]()
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
#endif
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK2
typedef __INTPTR_TYPE__ intptr_t;
// CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo();
template <typename T>
T tmain() {
#pragma omp target
#pragma omp teams distribute parallel for proc_bind(primary)
for(int i = 0; i < 1000; i++) {}
return T();
}
int main() {
// CK2-LABEL: @main
#pragma omp target
#pragma omp teams distribute parallel for proc_bind(primary)
for(int i = 0; i < 1000; i++) {}
return tmain<int>();
}
// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK2: call void [[OFFL1:@.+]]()
// CK2: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
// CK2: ret i32 [[CALL_RET]]
// CK2: define{{.+}} void [[OFFL1]](
// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK2: ret void
// CK2: define{{.+}} [[TMAIN]]()
// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK2: call void [[OFFL3:@.+]]()
// CK2: define{{.+}} [[OFFL2]]()
// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
// CK2: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK2: ret void
#endif
#endif #endif

clang/test/OpenMP/teams_distribute_parallel_for_simd_proc_bind_codegen.cpp

// add -fopenmp-targets // add -fopenmp-targets
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
// expected-no-diagnostics // expected-no-diagnostics
#ifndef HEADER #ifndef HEADER
#define HEADER #define HEADER
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1
// RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK1
typedef __INTPTR_TYPE__ intptr_t; typedef __INTPTR_TYPE__ intptr_t;
// CHECK-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* } // CK1-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" // CK1-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CHECK-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } // CK1-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo(); void foo();
struct S { struct S {
intptr_t a, b, c; intptr_t a, b, c;
S(intptr_t a) : a(a) {} S(intptr_t a) : a(a) {}
operator char() { return a; } operator char() { return a; }
~S() {} ~S() {}
}; };
template <typename T> template <typename T>
T tmain() { T tmain() {
#pragma omp target #pragma omp target
#pragma omp teams distribute parallel for simd proc_bind(master) #pragma omp teams distribute parallel for simd proc_bind(master)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
return T(); return T();
} }
int main() { int main() {
// CHECK-LABEL: @main // CK1-LABEL: @main
#pragma omp target #pragma omp target
#pragma omp teams distribute parallel for simd proc_bind(spread) #pragma omp teams distribute parallel for simd proc_bind(spread)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
#pragma omp target #pragma omp target
#pragma omp teams distribute parallel for simd proc_bind(close) #pragma omp teams distribute parallel for simd proc_bind(close)
for(int i = 0; i < 1000; i++) {} for(int i = 0; i < 1000; i++) {}
return tmain<int>(); return tmain<int>();
} }
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL1:@.+]]() // CK1: call void [[OFFL1:@.+]]()
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL2:@.+]]() // CK1: call void [[OFFL2:@.+]]()
// CHECK: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]() // CK1: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
// CHECK: ret i32 [[CALL_RET]] // CK1: ret i32 [[CALL_RET]]
// CHECK: define{{.+}} void [[OFFL1]]( // CK1: define{{.+}} void [[OFFL1]](
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 4)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
// CHECK: define{{.+}} [[OFFL2]]() // CK1: define{{.+}} [[OFFL2]]()
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 3)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
// CHECK: define{{.+}} [[TMAIN]]() // CK1: define{{.+}} [[TMAIN]]()
// CHECK: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}}) // CK1: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CHECK: call void [[OFFL3:@.+]]() // CK1: call void [[OFFL3:@.+]]()
// CHECK: define{{.+}} [[OFFL3]]() // CK1: define{{.+}} [[OFFL3]]()
// CHECK: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}}) // CK1: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
// CHECK: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]], // CK1: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
// CHECK: [[GTID_ADDR:%.+]] = alloca i32*, // CK1: [[GTID_ADDR:%.+]] = alloca i32*,
// CHECK: store i32* [[GTID_IN]], i32** [[GTID_ADDR]], // CK1: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CHECK: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]], // CK1: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CHECK: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]], // CK1: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CHECK: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2) // CK1: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 2)
// CHECK: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call( // CK1: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CHECK: ret void // CK1: ret void
// CK1: !{!"llvm.loop.vectorize.enable", i1 true}
#endif
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2
// RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
// RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK2
typedef __INTPTR_TYPE__ intptr_t;
// CHECK: !{!"llvm.loop.vectorize.enable", i1 true} // CK2-DAG: [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
// CK2-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
// CK2-DAG: [[DEF_LOC_2:@.+]] = private unnamed_addr constant [[IDENT_T_TY]] { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
void foo();
template <typename T>
T tmain() {
#pragma omp target
#pragma omp teams distribute parallel for simd proc_bind(primary)
for(int i = 0; i < 1000; i++) {}
return T();
}
int main() {
// CK2-LABEL: @main
#pragma omp target
#pragma omp teams distribute parallel for simd proc_bind(primary)
for(int i = 0; i < 1000; i++) {}
return tmain<int>();
}
// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK2: call void [[OFFL1:@.+]]()
// CK2: [[CALL_RET:%.+]] = call{{.+}} i32 [[TMAIN:@.+]]()
// CK2: ret i32 [[CALL_RET]]
// CK2: define{{.+}} void [[OFFL1]](
// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_1:@.+]] to {{.+}})
// CK2: define{{.+}} [[OMP_OUTLINED_1]](i32* {{.+}} [[GTID_IN:%.+]],
// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK2: ret void
// CK2: define{{.+}} [[TMAIN]]()
// CK2: call {{.*}}@__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, {{.+}})
// CK2: call void [[OFFL3:@.+]]()
// CK2: define{{.+}} [[OFFL2]]()
// CK2: call {{.*}}void {{.+}} @__kmpc_fork_teams({{.+}}, {{.+}}, {{.+}}* [[OMP_OUTLINED_3:@.+]] to {{.+}})
// CK2: define{{.+}} [[OMP_OUTLINED_3]](i32* {{.+}} [[GTID_IN:%.+]],
// CK2: [[GTID_ADDR:%.+]] = alloca i32*,
// CK2: store i32* [[GTID_IN]], i32** [[GTID_ADDR]],
// CK2: [[GTID_REF:%.+]] = load i32*, i32** [[GTID_ADDR]],
// CK2: [[GTID_VAL:%.+]] = load i32, i32* [[GTID_REF]],
// CK2: call {{.*}}void @__kmpc_push_proc_bind([[IDENT_T_TY]]* [[DEF_LOC_2]], i32 [[GTID_VAL]], i32 5)
// CK2: call {{.*}}void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(
// CK2: ret void
// CK2: !{!"llvm.loop.vectorize.enable", i1 true}
#endif
#endif #endif

llvm/include/llvm/Frontend/OpenMP/OMP.td

Show First 20 Lines • Show All 97 Lines • ▼ Show 20 Lines
} }
def OMPC_CopyPrivate : Clause<"copyprivate"> { def OMPC_CopyPrivate : Clause<"copyprivate"> {
let clangClass = "OMPCopyprivateClause"; let clangClass = "OMPCopyprivateClause";
let flangClass = "OmpObjectList"; let flangClass = "OmpObjectList";
} }
def OMP_PROC_BIND_master : ClauseVal<"master",2,1> {} def OMP_PROC_BIND_master : ClauseVal<"master",2,1> {}
def OMP_PROC_BIND_close : ClauseVal<"close",3,1> {} def OMP_PROC_BIND_close : ClauseVal<"close",3,1> {}
def OMP_PROC_BIND_spread : ClauseVal<"spread",4,1> {} def OMP_PROC_BIND_spread : ClauseVal<"spread",4,1> {}
def OMP_PROC_BIND_default : ClauseVal<"default",5,0> {} def OMP_PROC_BIND_primary : ClauseVal<"primary",5,1> {}
def OMP_PROC_BIND_unknown : ClauseVal<"unknown",6,0> { let isDefault = true; } def OMP_PROC_BIND_default : ClauseVal<"default",6,0> {}
def OMP_PROC_BIND_unknown : ClauseVal<"unknown",7,0> { let isDefault = true; }
def OMPC_ProcBind : Clause<"proc_bind"> { def OMPC_ProcBind : Clause<"proc_bind"> {
let clangClass = "OMPProcBindClause"; let clangClass = "OMPProcBindClause";
let flangClass = "OmpProcBindClause"; let flangClass = "OmpProcBindClause";
let enumClauseValue = "ProcBindKind"; let enumClauseValue = "ProcBindKind";
let allowedClauseValues = [ let allowedClauseValues = [
OMP_PROC_BIND_primary,
OMP_PROC_BIND_master, OMP_PROC_BIND_master,
OMP_PROC_BIND_close, OMP_PROC_BIND_close,
OMP_PROC_BIND_spread, OMP_PROC_BIND_spread,
OMP_PROC_BIND_default, OMP_PROC_BIND_default,
OMP_PROC_BIND_unknown OMP_PROC_BIND_unknown
]; ];
} }
▲ Show 20 Lines • Show All 1,540 Lines • Show Last 20 Lines

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

Show First 20 Lines • Show All 961 Lines • ▼ Show 20 Lines
#endif #endif
#define __OMP_PROC_BIND_KIND(Name, Value) \ #define __OMP_PROC_BIND_KIND(Name, Value) \
OMP_PROC_BIND_KIND(OMP_PROC_BIND_##Name, #Name, Value) OMP_PROC_BIND_KIND(OMP_PROC_BIND_##Name, #Name, Value)
__OMP_PROC_BIND_KIND(master, 2) __OMP_PROC_BIND_KIND(master, 2)
__OMP_PROC_BIND_KIND(close, 3) __OMP_PROC_BIND_KIND(close, 3)
__OMP_PROC_BIND_KIND(spread, 4) __OMP_PROC_BIND_KIND(spread, 4)
__OMP_PROC_BIND_KIND(primary, 5)
__OMP_PROC_BIND_KIND(default, 6) __OMP_PROC_BIND_KIND(default, 6)
__OMP_PROC_BIND_KIND(unknown, 7) __OMP_PROC_BIND_KIND(unknown, 7)
#undef __OMP_PROC_BIND_KIND #undef __OMP_PROC_BIND_KIND
#undef OMP_PROC_BIND_KIND #undef OMP_PROC_BIND_KIND
///} ///}
▲ Show 20 Lines • Show All 176 Lines • Show Last 20 Lines