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 -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 |
6 | // 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 |
7 | // 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 |
8 | // 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 |
9 | // 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 |
10 | // 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 |
11 | |
12 | // 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 |
13 | // 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 |
14 | // 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 |
15 | // 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 |
16 | // 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 |
17 | // 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 |
18 | |
19 | // Test target codegen - host bc file has to be created first. (no significant differences with host version of target region) |
20 | // 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 |
21 | // 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 |
22 | // 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 |
23 | // 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 |
24 | // 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 |
25 | // 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 |
26 | // 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 |
27 | // 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 |
28 | |
29 | // 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 |
30 | // 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 |
31 | // 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 |
32 | // 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 |
33 | // 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 |
34 | // 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 |
35 | // 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 |
36 | // 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 |
37 | // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} |
38 | |
39 | #ifdef CK1 |
40 | |
41 | // HCK1: define{{.*}} i32 @{{.+}}target_teams_fun{{.*}}( |
42 | int target_teams_fun(int *g){ |
43 | int n = 1000; |
44 | int a[1000]; |
45 | int te = n / 128; |
46 | int th = 128; |
47 | // discard n_addr |
48 | // HCK1: alloca i32, |
49 | // HCK1: [[TE:%.+]] alloca i32, |
50 | // HCK1: [[TH:%.+]] = alloca i32, |
51 | // HCK1: [[I:%.+]] = alloca i32, |
52 | // discard capture expressions for te and th |
53 | // HCK1: = alloca i32, |
54 | // HCK1: = alloca i32, |
55 | // HCK1: = alloca i32, |
56 | // HCK1: = alloca i32, |
57 | // HCK1: = alloca i32, |
58 | // HCK1: [[I_CAST:%.+]] = alloca i{{32|64}}, |
59 | // HCK1: [[N_CAST:%.+]] = alloca i{{32|64}}, |
60 | // HCK1: [[TE_CAST:%.+]] = alloca i{{32|64}}, |
61 | // HCK1: [[TH_CAST:%.+]] = alloca i{{32|64}}, |
62 | // HCK1: call void @__kmpc_push_target_tripcount(i64 -1, i64 %{{.+}}) |
63 | // HCK1: [[I_PAR:%.+]] = load{{.+}}, {{.+}} [[I_CAST]], |
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 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, |
68 | |
69 | // HCK1: call void @[[OFFL1:.+]](i{{32|64}} [[I_PAR]], i{{32|64}} [[N_PAR]], {{.+}}, i{{32|64}} [[TE_PAR]], i{{32|64}} [[TH_PAR]]) |
70 | int i; |
71 | #pragma omp target teams distribute parallel for simd num_teams(te), thread_limit(th) aligned(a : 8) safelen(16) simdlen(4) linear(i : n) |
72 | for(i = 0; i < n; i++) { |
73 | a[i] = 0; |
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 simd is_device_ptr(g) simdlen(8) |
80 | for( |
81 | int i = 0; i < n; i++) { |
82 | a[i] = g[0]; |
83 | } |
84 | }}} |
85 | |
86 | // outlined target regions |
87 | // HCK1: define internal void @[[OFFL1]](i{{32|64}} [[I_ARG:%.+]], i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]]) |
88 | // TCK1: define weak void @{{.+}}target_teams_fun{{.*}}(i{{32|64}} [[I_ARG:%.+]], i{{32|64}} [[N_ARG:%.+]], {{.+}}, i{{32|64}} [[TE_ARG:%.+]], i{{32|64}} [[TH_ARG:%.+]]) |
89 | // CK1: [[I_ADDR:%.+]] = alloca i{{32|64}}, |
90 | // CK1: [[N_ADDR:%.+]] = alloca i{{32|64}}, |
91 | // CK1: [[TE_ADDR:%.+]] = alloca i{{32|64}}, |
92 | // CK1: [[TH_ADDR:%.+]] = alloca i{{32|64}}, |
93 | // TCK1: store {{.+}} [[N_ARG]], {{.+}} [[N_ADDR]], |
94 | // CK1: store{{.+}} [[TE_ARG]], {{.+}} [[TE_ADDR]], |
95 | // CK1: store{{.+}} [[TH_ARG]], {{.+}} [[TH_ADDR]], |
96 | // CK1-64: [[TE_CONV:%.+]] = bitcast{{.+}} [[TE_ADDR]] to |
97 | // CK1-64: [[TH_CONV:%.+]] = bitcast{{.+}} [[TH_ADDR]] to |
98 | // CK1-64: [[TE_VAL:%.+]] = load i32, i32* [[TE_CONV]], |
99 | // CK1-64: [[TH_VAL:%.+]] = load i32, i32* [[TH_CONV]], |
100 | // CK1-32: [[TE_VAL:%.+]] = load i32, i32* [[TE_ADDR]], |
101 | // CK1-32: [[TH_VAL:%.+]] = load i32, i32* [[TH_ADDR]], |
102 | // CK1: {{%.+}} = call i32 @__kmpc_push_num_teams({{.+}}, {{.+}}, i32 [[TE_VAL]], i32 [[TH_VAL]]) |
103 | // CK1: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 3, {{.+}} @[[OUTL1:.+]] to {{.+}}, {{.+}}, {{.+}}) |
104 | // CK1: ret void |
105 | |
106 | // CK1: define internal void @[[OUTL1]]({{.+}}) |
107 | // CK1: [[ARRDECAY:%.+]] = getelementptr inbounds [1000 x i32], [1000 x i32]* %{{.+}}, i{{32|64}} 0, i{{32|64}} 0 |
108 | // CK1: [[ARR_CAST:%.+]] = ptrtoint i32* [[ARRDECAY]] to i{{32|64}} |
109 | // CK1: [[MASKED_PTR:%.+]] = and i{{32|64}} [[ARR_CAST]], 7 |
110 | // CK1: [[COND:%.+]] = icmp eq i{{32|64}} [[MASKED_PTR]], 0 |
111 | // CK1: call void @llvm.assume(i1 [[COND]]) |
112 | // CK1: call void @__kmpc_for_static_init_4( |
113 | // CK1: call void {{.+}} @__kmpc_fork_call( |
114 | // CK1: call void @__kmpc_for_static_fini( |
115 | // CK1: ret void |
116 | |
117 | // HCK1: define internal void @[[OFFL2]]( |
118 | // TCK1: define weak void @{{.+}}target_teams_fun{{.+}}( |
119 | // CK1: call void {{.+}} @__kmpc_fork_teams({{.+}}, i32 3, {{.+}} @[[OUTL2:.+]] to {{.+}}, {{.+}}, {{.+}}) |
120 | // CK1: ret void |
121 | |
122 | // CK1: define internal void @[[OUTL2]]({{.+}}) |
123 | // CK1: call void @__kmpc_for_static_init_4( |
124 | // CK1: call void {{.+}} @__kmpc_fork_call( |
125 | // CK1: call void @__kmpc_for_static_fini( |
126 | // CK1: ret void |
127 | |
128 | return a[0]; |
129 | } |
130 | |
131 | // CK1-DAG: !{!"llvm.loop.vectorize.width", i32 4} |
132 | // CK1-DAG: !{!"llvm.loop.vectorize.enable", i1 true} |
133 | // CK1-DAG: !{!"llvm.loop.vectorize.width", i32 8} |
134 | |
135 | #endif // CK1 |
136 | #endif // HEADER |
137 | |