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