1 | // expected-no-diagnostics |
2 | #ifndef HEADER |
3 | #define HEADER |
4 | // Test host codegen. |
5 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix HCK_NO_TGT |
6 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64 --check-prefix HCK1 --check-prefix HCK1-64 |
7 | // 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 |
8 | // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64 --check-prefix HCK1 --check-prefix HCK1-64 |
9 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 --check-prefix HCK1 --check-prefix HCK1-64 |
10 | // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s |
11 | // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 --check-prefix HCK1 --check-prefix HCK1-64 |
12 | |
13 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY1 |
14 | // 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 |
15 | // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY1 |
16 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix SIMD-ONLY1 |
17 | // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s |
18 | // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY1 |
19 | |
20 | // Test target codegen - host bc file has to be created first. (no significant differences with host version of target region) |
21 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc |
22 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -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 --check-prefix TCK1 --check-prefix TCK1-64 |
23 | // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s |
24 | // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64 --check-prefix TCK1 --check-prefix TCK1-64 |
25 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc |
26 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -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 --check-prefix TCK1 --check-prefix TCK1-32 |
27 | // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s |
28 | // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 --check-prefix TCK1 --check-prefix TCK1-32 |
29 | |
30 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc |
31 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix SIMD-ONLY1 |
32 | // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s |
33 | // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY1 |
34 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc |
35 | // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix SIMD-ONLY1 |
36 | // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s |
37 | // RUN: %clang_cc1 -DCK1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix SIMD-ONLY1 |
38 | // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} |
39 | |
40 | #ifdef CK1 |
41 | |
42 | // HCK_NO_TGT-NOT: @__kmpc_push_target_tripcount |
43 | |
44 | // HCK1: define{{.*}} i32 @{{.+}}target_teams_fun{{.*}}( |
45 | int target_teams_fun(int *g){ |
46 | int n = 1000; |
47 | int a[1000]; |
48 | int te = n / 128; |
49 | int th = 128; |
50 | // discard n_addr |
51 | // HCK1: alloca i32, |
52 | // HCK1: [[TE:%.+]] alloca i32, |
53 | // HCK1: [[TH:%.+]] = alloca i32, |
54 | // discard capture expressions for te and th |
55 | // HCK1: = alloca i32, |
56 | // HCK1: = alloca i32, |
57 | // HCK1: = alloca i32, |
58 | // HCK1: = alloca i32, |
59 | // HCK1: = alloca i32, |
60 | // HCK1: [[N_CAST:%.+]] = alloca i{{32|64}}, |
61 | // HCK1: [[TE_CAST:%.+]] = alloca i{{32|64}}, |
62 | // HCK1: [[TH_CAST:%.+]] = alloca i{{32|64}}, |
63 | // HCK1: call void @__kmpc_push_target_tripcount(i64 -1, i64 %{{.+}}) |
64 | // HCK1: [[N_PAR:%.+]] = load{{.+}}, {{.+}} [[N_CAST]], |
65 | // HCK1: [[TE_PAR:%.+]] = load{{.+}}, {{.+}} [[TE_CAST]], |
66 | // HCK1: [[TH_PAR:%.+]] = load{{.+}}, {{.+}} [[TH_CAST]], |
67 | // HCK1: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 4, i8** %{{[^,]+}}, i8** %{{[^,]+}}, |
68 | |
69 | // HCK1: call void @[[OFFL1:.+]](i{{32|64}} [[N_PAR]], {{.+}}, i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]]) |
70 | #pragma omp target teams distribute parallel for num_teams(te), thread_limit(th) |
71 | for(int i = 0; i < n; i++) { |
72 | a[i] = 0; |
73 | #pragma omp cancel for |
74 | } |
75 | |
76 | // HCK1: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), |
77 | // HCK1: call void @[[OFFL2:.+]](i{{64|32}} %{{.+}}) |
78 | {{{ |
79 | #pragma omp target teams distribute parallel for is_device_ptr(g) |
80 | for(int i = 0; i < n; i++) { |
81 | a[i] = g[0]; |
82 | } |
83 | }}} |
84 | |
85 | // outlined target regions |
86 | // HCK1: define internal void @[[OFFL1]](i{{32|64}} [[N_ARG:%.+]], i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]]) |
87 | // TCK1: define weak void @{{.+}}target_teams_fun{{.*}}(i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]]) |
88 | // CK1: [[N_ADDR:%.+]] = alloca i{{32|64}}, |
89 | // CK1: [[TE_ADDR:%.+]] = alloca i{{32|64}}, |
90 | // CK1: [[TH_ADDR:%.+]] = alloca i{{32|64}}, |
91 | // TCK1: store {{.+}} [[N_ARG]], {{.+}} [[N_ADDR]], |
92 | // CK1: store{{.+}} [[TE_ARG]], {{.+}} [[TE_ADDR]], |
93 | // CK1: store{{.+}} [[TH_ARG]], {{.+}} [[TH_ADDR]], |
94 | // CK1-64: [[TE_CONV:%.+]] = bitcast{{.+}} [[TE_ADDR]] to |
95 | // CK1-64: [[TH_CONV:%.+]] = bitcast{{.+}} [[TH_ADDR]] to |
96 | // CK1-64: [[TE_VAL:%.+]] = load i32, i32* [[TE_CONV]], |
97 | // CK1-64: [[TH_VAL:%.+]] = load i32, i32* [[TH_CONV]], |
98 | // CK1-32: [[TE_VAL:%.+]] = load i32, i32* [[TE_ADDR]], |
99 | // CK1-32: [[TH_VAL:%.+]] = load i32, i32* [[TH_ADDR]], |
100 | // CK1: {{%.+}} = call i32 @__kmpc_push_num_teams({{.+}}, {{.+}}, i32 [[TE_VAL]], i32 [[TH_VAL]]) |
101 | // CK1: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 2, {{.+}} @[[OUTL1:.+]] to {{.+}}, {{.+}}, {{.+}}) |
102 | // CK1: ret void |
103 | |
104 | // CK1: define internal void @[[OUTL1]]({{.+}}) |
105 | // CK1: call void @__kmpc_for_static_init_4( |
106 | // CK1: call void {{.+}} @__kmpc_fork_call( |
107 | // CK1: call void @__kmpc_for_static_fini( |
108 | // CK1: ret void |
109 | |
110 | // HCK1: define internal void @[[OFFL2]]( |
111 | // TCK1: define weak void @{{.+}}target_teams_fun{{.+}}( |
112 | // CK1: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 3, {{.+}} @[[OUTL2:.+]] to {{.+}}, {{.+}}, {{.+}}) |
113 | // CK1: ret void |
114 | |
115 | // CK1: define internal void @[[OUTL2]]({{.+}}) |
116 | // CK1: call void @__kmpc_for_static_init_4( |
117 | // CK1: call void {{.+}} @__kmpc_fork_call( |
118 | // CK1: call void @__kmpc_for_static_fini( |
119 | // CK1: ret void |
120 | |
121 | return a[0]; |
122 | } |
123 | |
124 | #endif // CK1 |
125 | #endif // HEADER |
126 | |