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 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 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 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 11 #ifdef CK1 12 13 int Gbla; 14 long long Gblb; 15 int &Gblc = Gbla; 16 17 // CK1-LABEL: teams_argument_global_local 18 int teams_argument_global_local(int a){ 19 int comp = 1; 20 21 int la = 23; 22 float lc = 25.0; 23 24 // CK1: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) 25 // CK1: call void @{{.+}}(i{{64|32}} %{{.+}}) 26 #pragma omp target 27 #pragma omp teams 28 { 29 ++comp; 30 } 31 32 // CK1: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0) 33 // CK1: call void @{{.+}}(i{{64|32}} %{{.+}}) 34 #pragma omp target 35 {{{ 36 #pragma omp teams 37 { 38 ++comp; 39 } 40 }}} 41 42 // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 0) 43 // CK1-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]], 44 45 // CK1: call void @{{.+}}(i{{64|32}} %{{.+}}) 46 #pragma omp target 47 #pragma omp teams num_teams(la) 48 { 49 ++comp; 50 } 51 52 // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 [[NT:%[^,]+]]) 53 // CK1-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]], 54 55 // CK1: call void @{{.+}}(i{{64|32}} %{{.+}}) 56 #pragma omp target 57 #pragma omp teams thread_limit(la) 58 { 59 ++comp; 60 } 61 62 // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 5, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]]) 63 64 // CK1-DAG: [[NT]] = add nsw i32 [[NTA:%[^,]+]], [[NTB:%[^,]+]] 65 // CK1-DAG: [[NTA]] = load i32, i32* @Gbla, 66 // CK1-DAG: [[NTB]] = load i32, i32* %{{.+}}, 67 68 // CK1-DAG: [[TL]] = trunc i64 [[TLA:%[^,]+]] to i32 69 // CK1-DAG: [[TLA]] = add nsw i64 [[TLB:%[^,]+]], [[TLC:%[^,]+]] 70 // CK1-DAG: [[TLC]] = fptosi float [[TLD:%[^,]+]] to i64 71 // CK1-DAG: [[TLD]] = load float, float* %{{.+}}, 72 // CK1-DAG: [[TLB]] = load i64, i64* @Gblb, 73 74 // CK1: call void @{{.+}}(i{{.+}} {{.+}}, i{{.+}} {{.+}}, i{{.+}} {{.+}}, i{{.+}} {{.+}}, i{{.+}} {{.+}}) 75 #pragma omp target 76 #pragma omp teams num_teams(Gbla+a) thread_limit(Gblb+(long long)lc) 77 { 78 ++comp; 79 } 80 81 // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 {{.+}}, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]]) 82 83 // CK1-DAG: [[NT]] = add nsw i32 [[NTA:%[^,]+]], 1 84 // CK1-DAG: [[NTA]] = load i32, i32* @Gbla, 85 86 // CK1-DAG: [[TL]] = add nsw i32 [[TLA:%[^,]+]], 2 87 // CK1-DAG: [[TLA]] = load i32, i32* @Gbla, 88 89 // CK1: call void @{{.+}}(i{{.+}} {{.+}} 90 #pragma omp target 91 #pragma omp teams num_teams(Gblc+1) thread_limit(Gblc+2) 92 { 93 comp += Gblc; 94 } 95 96 return comp; 97 } 98 99 #endif // CK1 100 101 // Test host codegen. 102 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64 103 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 104 // RUN: %clang_cc1 -DCK2 -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 CK2 --check-prefix CK2-64 105 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32 106 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 107 // RUN: %clang_cc1 -DCK2 -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 CK2 --check-prefix CK2-32 108 #ifdef CK2 109 110 // CK2-DAG: [[SSI:%.+]] = type { i32, float } 111 // CK2-DAG: [[SSL:%.+]] = type { i64, float } 112 template <typename T> 113 struct SS{ 114 T a; 115 float b; 116 }; 117 118 SS<int> Gbla; 119 SS<long long> Gblb; 120 121 // CK2-LABEL: teams_template_arg 122 int teams_template_arg(void) { 123 int comp = 1; 124 125 SS<int> la; 126 SS<long long> lb; 127 128 // CK2-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]]) 129 130 // CK2-DAG: [[NT]] = load i32, i32* getelementptr inbounds ([[SSI]], [[SSI]]* @Gbla, i32 0, i32 0) 131 132 // CK2-DAG: [[TL]] = trunc i64 [[TLA:%[^,]+]] to i32 133 // CK2-DAG: [[TLA]] = fptosi float [[TLB:%[^,]+]] to i64 134 // CK2-DAG: [[TLB]] = load float, float* [[TLC:%[^,]+]], 135 // CK2-DAG: [[TLC]] = getelementptr inbounds [[SSI]], [[SSI]]* %{{.+}}, i32 0, i32 1 136 137 // CK2: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}}, {{.+}} {{.+}}) 138 #pragma omp target 139 #pragma omp teams num_teams(Gbla.a) thread_limit((long long)la.b) 140 { 141 ++comp; 142 } 143 144 // CK2-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 [[TL:%[^,]+]]) 145 146 // CK2-DAG: [[TL]] = trunc i64 [[TLD:%[^,]+]] to i32 147 // CK2-DAG: [[TLD]] = load i64, i64* getelementptr inbounds ([[SSL]], [[SSL]]* @Gblb, i32 0, i32 0), 148 149 // CK2-DAG: [[NT]] = trunc i64 [[NTA:%[^,]+]] to i32 150 // CK2-DAG: [[NTA]] = fptosi float [[NTB:%[^,]+]] to i64 151 // CK2-DAG: [[NTB]] = load float, float* [[NTC:%[^,]+]], 152 // CK2-DAG: [[NTC]] = getelementptr inbounds [[SSL]], [[SSL]]* %{{.+}}, i32 0, i32 1 153 154 // CK2: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}}, {{.+}} {{.+}}) 155 #pragma omp target 156 #pragma omp teams num_teams((long long)lb.b) thread_limit(Gblb.a) 157 { 158 ++comp; 159 } 160 return comp; 161 } 162 #endif // CK2 163 164 // Test host codegen. 165 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64 166 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s 167 // RUN: %clang_cc1 -DCK3 -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 CK3 --check-prefix CK3-64 168 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32 169 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s 170 // RUN: %clang_cc1 -DCK3 -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 CK3 --check-prefix CK3-32 171 #ifdef CK3 172 173 // CK3: [[SSI:%.+]] = type { i32, float } 174 // CK3-LABEL: teams_template_struct 175 176 template <typename T, int X, long long Y> 177 struct SS{ 178 T a; 179 float b; 180 181 int foo(void) { 182 int comp = 1; 183 184 // CK3-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 123) 185 186 // CK3-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]], 187 // CK3-DAG: [[NTA]] = getelementptr inbounds [[SSI]], [[SSI]]* [[NTB:%[^,]+]], i32 0, i32 0 188 // CK3-DAG: [[NTB]] = load [[SSI]]*, [[SSI]]** %{{.+}}, 189 190 // CK3: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}}) 191 #pragma omp target 192 #pragma omp teams num_teams(a) thread_limit(X) 193 { 194 ++comp; 195 } 196 197 // CK3-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 456, i32 [[TL:%[^,]+]]) 198 199 // CK3-DAG: [[TL]] = add nsw i32 [[TLA:%[^,]+]], 123 200 // CK3-DAG: [[TLA]] = fptosi float [[TLB:%[^,]+]] to i32 201 // CK3-DAG: [[TLB]] = load float, float* [[TLC:%[^,]+]], 202 // CK3-DAG: [[TLC]] = getelementptr inbounds [[SSI]], [[SSI]]* [[THIS:%[^,]+]], i32 0, i32 1 203 204 // CK3: call void @{{.+}}({{.+}} {{.+}}, {{.+}} {{.+}}) 205 #pragma omp target 206 #pragma omp teams num_teams(Y) thread_limit((int)b+X) 207 { 208 ++comp; 209 } 210 return comp; 211 } 212 }; 213 214 int teams_template_struct(void) { 215 SS<int, 123, 456> V; 216 return V.foo(); 217 218 } 219 #endif // CK3 220 221 // Test target codegen - host bc file has to be created first. 222 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 223 // RUN: %clang_cc1 -DCK4 -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 CK4 --check-prefix CK4-64 224 // RUN: %clang_cc1 -DCK4 -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 225 // RUN: %clang_cc1 -DCK4 -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 CK4 --check-prefix CK4-64 226 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 227 // RUN: %clang_cc1 -DCK4 -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 CK4 --check-prefix CK4-32 228 // RUN: %clang_cc1 -DCK4 -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 229 // RUN: %clang_cc1 -DCK4 -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 CK4 --check-prefix CK4-32 230 231 #ifdef CK4 232 233 // CK4-DAG: %ident_t = type { i32, i32, i32, i32, i8* } 234 // CK4-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" 235 // CK4-DAG: [[DEF_LOC_0:@.+]] = private unnamed_addr constant %ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } 236 // CK4-DEBUG-DAG: [[LOC1:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;main;[[@LINE+14]];9;;\00" 237 // CK4-DEBUG-DAG: [[LOC2:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;tmain;[[@LINE+7]];9;;\00" 238 239 template <typename T> 240 int tmain(T argc) { 241 #pragma omp target 242 #pragma omp teams 243 argc = 0; 244 return 0; 245 } 246 247 int main (int argc, char **argv) { 248 #pragma omp target 249 #pragma omp teams 250 argc = 0; 251 return tmain(argv); 252 } 253 254 // CK4: define {{.*}}void @{{[^,]+}}(i{{.+}} %[[ARGC:.+]]) 255 // CK4: [[ARGCADDR:%.+]] = alloca i{{.+}} 256 // CK4: store i{{.+}} %[[ARGC]], i{{.+}}* [[ARGCADDR]] 257 // CK4-64: [[CONV:%.+]] = bitcast i64* [[ARGCADDR]] to i32* 258 // CK4-64: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* {{.+}} to void (i32*, i32*, ...)*), i32* [[CONV]]) 259 // CK4-32: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* {{.+}} to void (i32*, i32*, ...)*), i32* [[ARGCADDR]]) 260 // CK4: ret void 261 // CK4-NEXT: } 262 263 // CK4: define {{.*}}void @{{[^,]+}}(i8** [[ARGC1:%.+]]) 264 // CK4: [[ARGCADDR1:%.+]] = alloca i8** 265 // CK4: store i8** [[ARGC1]], i8*** [[ARGCADDR1]] 266 // CK4: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* {{.+}} to void (i32*, i32*, ...)*), i8*** [[ARGCADDR1]]) 267 268 269 #endif // CK4 270 271 // Test target codegen - host bc file has to be created first. 272 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc 273 // RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-64 274 // RUN: %clang_cc1 -DCK5 -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 275 // RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-64 276 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc 277 // RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-32 278 // RUN: %clang_cc1 -DCK5 -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 279 // RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-32 280 281 // expected-no-diagnostics 282 #ifdef CK5 283 284 // CK5-DAG: %ident_t = type { i32, i32, i32, i32, i8* } 285 // CK5-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00" 286 // CK5-DAG: [[DEF_LOC_0:@.+]] = private unnamed_addr constant %ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) } 287 // CK5-DEBUG-DAG: [[LOC1:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;main;[[@LINE+14]];9;;\00" 288 // CK5-DEBUG-DAG: [[LOC2:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}teams_codegen.cpp;tmain;[[@LINE+7]];9;;\00" 289 290 template <typename T> 291 int tmain(T argc) { 292 int a = 10; 293 int b = 5; 294 #pragma omp target 295 #pragma omp teams num_teams(a) thread_limit(b) 296 { 297 argc = 0; 298 } 299 return 0; 300 } 301 302 int main (int argc, char **argv) { 303 int a = 20; 304 int b = 5; 305 #pragma omp target 306 #pragma omp teams num_teams(a) thread_limit(b) 307 { 308 argc = 0; 309 } 310 return tmain(argv); 311 } 312 313 // CK5: define {{.*}}void @{{[^,]+}}(i{{.+}} [[AP:%.+]], i{{.+}} [[BP:%.+]], i{{.+}} [[ARGC:.+]]) 314 // CK5: [[AADDR:%.+]] = alloca i{{.+}} 315 // CK5: [[BADDR:%.+]] = alloca i{{.+}} 316 // CK5: [[ARGCADDR:%.+]] = alloca i{{.+}} 317 // CK5: [[GBL_TH_NUM:%.+]] = call i32 @__kmpc_global_thread_num(%ident_t* [[DEF_LOC_0]]) 318 // CK5: store i{{.+}} [[AP]], i{{.+}}* [[AADDR]] 319 // CK5: store i{{.+}} [[BP]], i{{.+}}* [[BADDR]] 320 // CK5: store i{{.+}} [[ARGC]], i{{.+}}* [[ARGCADDR]] 321 // CK5-64: [[ACONV:%.+]] = bitcast i64* [[AADDR]] to i32* 322 // CK5-64: [[BCONV:%.+]] = bitcast i64* [[BADDR]] to i32* 323 // CK5-64: [[CONV:%.+]] = bitcast i64* [[ARGCADDR]] to i32* 324 // CK5-64: [[ACONVVAL:%.+]] = load i32, i32* [[ACONV]] 325 // CK5-64: [[BCONVVAL:%.+]] = load i32, i32* [[BCONV]] 326 // CK5-32: [[ACONVVAL:%.+]] = load i32, i32* [[AADDR]] 327 // CK5-32: [[BCONVVAL:%.+]] = load i32, i32* [[BADDR]] 328 // CK5: {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[ACONVVAL]], i32 [[BCONVVAL]]) 329 // CK5-64: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[CONV]]) 330 // CK5-32: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[ARGCADDR]]) 331 332 // CK5: define {{.*}}void @{{[^,]+}}(i{{.+}} [[AP:%.+]], i{{.+}} [[BP:%.+]], i{{.+}}** [[ARGC:%.+]]) 333 // CK5: [[AADDR:%.+]] = alloca i{{.+}} 334 // CK5: [[BADDR:%.+]] = alloca i{{.+}} 335 // CK5: [[ARGCADDR:%.+]] = alloca i{{.+}}** 336 // CK5: [[GBL_TH_NUM:%.+]] = call i32 @__kmpc_global_thread_num(%ident_t* [[DEF_LOC_0]]) 337 // CK5: store i{{.+}} [[AP]], i{{.+}}* [[AADDR]] 338 // CK5: store i{{.+}} [[BP]], i{{.+}}* [[BADDR]] 339 // CK5: store i{{.+}}** [[ARGC]], i{{.+}}*** [[ARGCADDR]] 340 // CK5-64: [[ACONV:%.+]] = bitcast i64* [[AADDR]] to i32* 341 // CK5-64: [[BCONV:%.+]] = bitcast i64* [[BADDR]] to i32* 342 // CK5-64: [[ACONVVAL:%.+]] = load i32, i32* [[ACONV]] 343 // CK5-64: [[BCONVVAL:%.+]] = load i32, i32* [[BCONV]] 344 // CK5-64: {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[ACONVVAL]], i32 [[BCONVVAL]]) 345 // CK5-32: [[A_VAL:%.+]] = load i32, i32* [[AADDR]] 346 // CK5-32: [[B_VAL:%.+]] = load i32, i32* [[BADDR]] 347 // CK5-32: {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[A_VAL]], i32 [[B_VAL]]) 348 // CK5: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i{{.+}})* @.omp_outlined.{{.+}} to void (i32*, i32*, ...)*), i{{.+}}*** [[ARGCADDR]]) 349 // CK5: ret void 350 // CK5-NEXT: } 351 352 #endif // CK5 353 #endif 354
