1 // Test target codegen - host bc file has to be created first. 2 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc 3 // RUN: %clang_cc1 -verify -fopenmp -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 4 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc 5 // RUN: %clang_cc1 -verify -fopenmp -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 6 // RUN: %clang_cc1 -verify -fopenmp -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 7 // expected-no-diagnostics 8 #ifndef HEADER 9 #define HEADER 10 11 // CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1 12 // CHECK: @__omp_offloading_{{.+}}_l52_exec_mode = weak constant i8 1 13 // CHECK-NOT: @__omp_offloading_{{.+}}_exec_mode = weak constant i8 1 14 15 void foo() { 16 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 17 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 18 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 19 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 20 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 21 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 22 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 23 #pragma omp target teams distribute parallel for simd 24 for (int i = 0; i < 10; ++i) 25 ; 26 #pragma omp target teams distribute parallel for simd schedule(static) 27 for (int i = 0; i < 10; ++i) 28 ; 29 #pragma omp target teams distribute parallel for simd schedule(static, 1) 30 for (int i = 0; i < 10; ++i) 31 ; 32 #pragma omp target teams distribute parallel for simd schedule(auto) 33 for (int i = 0; i < 10; ++i) 34 ; 35 #pragma omp target teams distribute parallel for simd schedule(runtime) 36 for (int i = 0; i < 10; ++i) 37 ; 38 #pragma omp target teams distribute parallel for simd schedule(dynamic) 39 for (int i = 0; i < 10; ++i) 40 ; 41 #pragma omp target teams distribute parallel for simd schedule(guided) 42 for (int i = 0; i < 10; ++i) 43 ; 44 int a; 45 // CHECK: call void @__kmpc_kernel_init( 46 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 47 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 48 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 49 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 50 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 51 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 52 #pragma omp target teams distribute parallel for lastprivate(a) 53 for (int i = 0; i < 10; ++i) 54 a = i; 55 #pragma omp target teams distribute parallel for schedule(static) 56 for (int i = 0; i < 10; ++i) 57 ; 58 #pragma omp target teams distribute parallel for schedule(static, 1) 59 for (int i = 0; i < 10; ++i) 60 ; 61 #pragma omp target teams distribute parallel for schedule(auto) 62 for (int i = 0; i < 10; ++i) 63 ; 64 #pragma omp target teams distribute parallel for schedule(runtime) 65 for (int i = 0; i < 10; ++i) 66 ; 67 #pragma omp target teams distribute parallel for schedule(dynamic) 68 for (int i = 0; i < 10; ++i) 69 ; 70 #pragma omp target teams distribute parallel for schedule(guided) 71 for (int i = 0; i < 10; ++i) 72 ; 73 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 74 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 75 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 76 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 77 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 78 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 79 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 80 #pragma omp target teams 81 #pragma omp distribute parallel for simd 82 for (int i = 0; i < 10; ++i) 83 ; 84 #pragma omp target teams 85 #pragma omp distribute parallel for simd schedule(static) 86 for (int i = 0; i < 10; ++i) 87 ; 88 #pragma omp target teams 89 #pragma omp distribute parallel for simd schedule(static, 1) 90 for (int i = 0; i < 10; ++i) 91 ; 92 #pragma omp target teams 93 #pragma omp distribute parallel for simd schedule(auto) 94 for (int i = 0; i < 10; ++i) 95 ; 96 #pragma omp target teams 97 #pragma omp distribute parallel for simd schedule(runtime) 98 for (int i = 0; i < 10; ++i) 99 ; 100 #pragma omp target teams 101 #pragma omp distribute parallel for simd schedule(dynamic) 102 for (int i = 0; i < 10; ++i) 103 ; 104 #pragma omp target teams 105 #pragma omp distribute parallel for simd schedule(guided) 106 for (int i = 0; i < 10; ++i) 107 ; 108 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 109 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 110 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 111 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 112 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 113 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 114 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 115 #pragma omp target teams 116 #pragma omp distribute parallel for 117 for (int i = 0; i < 10; ++i) 118 ; 119 #pragma omp target teams 120 #pragma omp distribute parallel for schedule(static) 121 for (int i = 0; i < 10; ++i) 122 ; 123 #pragma omp target teams 124 #pragma omp distribute parallel for schedule(static, 1) 125 for (int i = 0; i < 10; ++i) 126 ; 127 #pragma omp target teams 128 #pragma omp distribute parallel for schedule(auto) 129 for (int i = 0; i < 10; ++i) 130 ; 131 #pragma omp target teams 132 #pragma omp distribute parallel for schedule(runtime) 133 for (int i = 0; i < 10; ++i) 134 ; 135 #pragma omp target teams 136 #pragma omp distribute parallel for schedule(dynamic) 137 for (int i = 0; i < 10; ++i) 138 ; 139 #pragma omp target teams 140 #pragma omp distribute parallel for schedule(guided) 141 for (int i = 0; i < 10; ++i) 142 ; 143 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 144 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 145 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 146 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 147 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 148 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 149 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 150 #pragma omp target 151 #pragma omp teams 152 #pragma omp distribute parallel for 153 for (int i = 0; i < 10; ++i) 154 ; 155 #pragma omp target 156 #pragma omp teams 157 #pragma omp distribute parallel for schedule(static) 158 for (int i = 0; i < 10; ++i) 159 ; 160 #pragma omp target 161 #pragma omp teams 162 #pragma omp distribute parallel for schedule(static, 1) 163 for (int i = 0; i < 10; ++i) 164 ; 165 #pragma omp target 166 #pragma omp teams 167 #pragma omp distribute parallel for schedule(auto) 168 for (int i = 0; i < 10; ++i) 169 ; 170 #pragma omp target 171 #pragma omp teams 172 #pragma omp distribute parallel for schedule(runtime) 173 for (int i = 0; i < 10; ++i) 174 ; 175 #pragma omp target 176 #pragma omp teams 177 #pragma omp distribute parallel for schedule(dynamic) 178 for (int i = 0; i < 10; ++i) 179 ; 180 #pragma omp target 181 #pragma omp teams 182 #pragma omp distribute parallel for schedule(guided) 183 for (int i = 0; i < 10; ++i) 184 ; 185 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 186 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 187 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 188 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 189 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 190 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 191 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 192 #pragma omp target parallel for 193 for (int i = 0; i < 10; ++i) 194 ; 195 #pragma omp target parallel for schedule(static) 196 for (int i = 0; i < 10; ++i) 197 ; 198 #pragma omp target parallel for schedule(static, 1) 199 for (int i = 0; i < 10; ++i) 200 ; 201 #pragma omp target parallel for schedule(auto) 202 for (int i = 0; i < 10; ++i) 203 ; 204 #pragma omp target parallel for schedule(runtime) 205 for (int i = 0; i < 10; ++i) 206 ; 207 #pragma omp target parallel for schedule(dynamic) 208 for (int i = 0; i < 10; ++i) 209 ; 210 #pragma omp target parallel for schedule(guided) 211 for (int i = 0; i < 10; ++i) 212 ; 213 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 214 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 215 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 216 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 217 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 218 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 219 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 220 #pragma omp target parallel 221 #pragma omp for simd 222 for (int i = 0; i < 10; ++i) 223 ; 224 #pragma omp target parallel 225 #pragma omp for simd schedule(static) 226 for (int i = 0; i < 10; ++i) 227 ; 228 #pragma omp target parallel 229 #pragma omp for simd schedule(static, 1) 230 for (int i = 0; i < 10; ++i) 231 ; 232 #pragma omp target parallel 233 #pragma omp for simd schedule(auto) 234 for (int i = 0; i < 10; ++i) 235 ; 236 #pragma omp target parallel 237 #pragma omp for simd schedule(runtime) 238 for (int i = 0; i < 10; ++i) 239 ; 240 #pragma omp target parallel 241 #pragma omp for simd schedule(dynamic) 242 for (int i = 0; i < 10; ++i) 243 ; 244 #pragma omp target parallel 245 #pragma omp for simd schedule(guided) 246 for (int i = 0; i < 10; ++i) 247 ; 248 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 249 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 250 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 251 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 252 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 253 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 254 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 255 #pragma omp target 256 #pragma omp parallel 257 #pragma omp for simd ordered 258 for (int i = 0; i < 10; ++i) 259 ; 260 #pragma omp target 261 #pragma omp parallel 262 #pragma omp for simd schedule(static) 263 for (int i = 0; i < 10; ++i) 264 ; 265 #pragma omp target 266 #pragma omp parallel 267 #pragma omp for simd schedule(static, 1) 268 for (int i = 0; i < 10; ++i) 269 ; 270 #pragma omp target 271 #pragma omp parallel 272 #pragma omp for simd schedule(auto) 273 for (int i = 0; i < 10; ++i) 274 ; 275 #pragma omp target 276 #pragma omp parallel 277 #pragma omp for simd schedule(runtime) 278 for (int i = 0; i < 10; ++i) 279 ; 280 #pragma omp target 281 #pragma omp parallel 282 #pragma omp for simd schedule(dynamic) 283 for (int i = 0; i < 10; ++i) 284 ; 285 #pragma omp target 286 #pragma omp parallel 287 #pragma omp for simd schedule(guided) 288 for (int i = 0; i < 10; ++i) 289 ; 290 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 291 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 292 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 0, i16 0) 293 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 294 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 295 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 296 // CHECK: call void @__kmpc_spmd_kernel_init(i32 {{.+}}, i16 1, i16 {{.+}}) 297 #pragma omp target 298 #pragma omp parallel for 299 for (int i = 0; i < 10; ++i) 300 ; 301 #pragma omp target 302 #pragma omp parallel for schedule(static) 303 for (int i = 0; i < 10; ++i) 304 ; 305 #pragma omp target 306 #pragma omp parallel for schedule(static, 1) 307 for (int i = 0; i < 10; ++i) 308 ; 309 #pragma omp target 310 #pragma omp parallel for schedule(auto) 311 for (int i = 0; i < 10; ++i) 312 ; 313 #pragma omp target 314 #pragma omp parallel for schedule(runtime) 315 for (int i = 0; i < 10; ++i) 316 ; 317 #pragma omp target 318 #pragma omp parallel for schedule(dynamic) 319 for (int i = 0; i < 10; ++i) 320 ; 321 #pragma omp target 322 #pragma omp parallel for schedule(guided) 323 for (int i = 0; i < 10; ++i) 324 ; 325 } 326 327 #endif 328 329