1243 lines · cpp
1// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s2 3extern "C" void acc_combined(int N, int cond) {4 // CHECK: cir.func{{.*}} @acc_combined(%[[ARG_N:.*]]: !s32i loc{{.*}}, %[[ARG_COND:.*]]: !s32i loc{{.*}}) {5 // CHECK-NEXT: %[[ALLOCA_N:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["N", init]6 // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]7 // CHECK-NEXT: cir.store %[[ARG_N]], %[[ALLOCA_N]] : !s32i, !cir.ptr<!s32i>8 // CHECK-NEXT: cir.store %[[ARG_COND]], %[[COND]] : !s32i, !cir.ptr<!s32i>9 10#pragma acc parallel loop11 for(unsigned I = 0; I < N; ++I);12 // CHECK: acc.parallel combined(loop) {13 // CHECK: acc.loop combined(parallel) {14 // CHECK: acc.yield15 // CHECK-NEXT: } loc16 // CHECK: acc.yield17 // CHECK-NEXT: } loc18 19#pragma acc serial loop20 for(unsigned I = 0; I < N; ++I);21 // CHECK: acc.serial combined(loop) {22 // CHECK: acc.loop combined(serial) {23 // CHECK: acc.yield24 // CHECK-NEXT: } loc25 // CHECK: acc.yield26 // CHECK-NEXT: } loc27 28#pragma acc kernels loop29 for(unsigned I = 0; I < N; ++I);30 // CHECK: acc.kernels combined(loop) {31 // CHECK: acc.loop combined(kernels) {32 // CHECK: acc.yield33 // CHECK-NEXT: } loc34 // CHECK: acc.terminator35 // CHECK-NEXT: } loc36 37#pragma acc parallel loop default(none)38 for(unsigned I = 0; I < N; ++I);39 // CHECK: acc.parallel combined(loop) {40 // CHECK: acc.loop combined(parallel) {41 // CHECK: acc.yield42 // CHECK-NEXT: } loc43 // CHECK: acc.yield44 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>} loc45 46#pragma acc serial loop default(present)47 for(unsigned I = 0; I < N; ++I);48 // CHECK: acc.serial combined(loop) {49 // CHECK: acc.loop combined(serial) {50 // CHECK: acc.yield51 // CHECK-NEXT: } loc52 // CHECK: acc.yield53 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>} loc54 55#pragma acc kernels loop default(none)56 for(unsigned I = 0; I < N; ++I);57 // CHECK: acc.kernels combined(loop) {58 // CHECK: acc.loop combined(kernels) {59 // CHECK: acc.yield60 // CHECK-NEXT: } loc61 // CHECK: acc.terminator62 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>} loc63 64#pragma acc parallel loop seq65 for(unsigned I = 0; I < N; ++I);66 // CHECK: acc.parallel combined(loop) {67 // CHECK: acc.loop combined(parallel) {68 // CHECK: acc.yield69 // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc70 // CHECK: acc.yield71 // CHECK-NEXT: } loc72#pragma acc serial loop device_type(nvidia, radeon) seq73 for(unsigned I = 0; I < N; ++I);74 // CHECK: acc.serial combined(loop) {75 // CHECK: acc.loop combined(serial) {76 // CHECK: acc.yield77 // CHECK-NEXT: } attributes {seq = [#acc.device_type<nvidia>, #acc.device_type<radeon>, #acc.device_type<none>]} loc78 // CHECK: acc.yield79 // CHECK-NEXT: } loc80#pragma acc kernels loop seq device_type(nvidia, radeon)81 for(unsigned I = 0; I < N; ++I);82 // CHECK: acc.kernels combined(loop) {83 // CHECK: acc.loop combined(kernels) {84 // CHECK: acc.yield85 // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc86 // CHECK: acc.terminator87 // CHECK-NEXT: } loc88 89#pragma acc parallel loop auto90 for(unsigned I = 0; I < N; ++I);91 // CHECK: acc.parallel combined(loop) {92 // CHECK: acc.loop combined(parallel) {93 // CHECK: acc.yield94 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc95 // CHECK: acc.yield96 // CHECK-NEXT: } loc97#pragma acc serial loop device_type(nvidia, radeon) auto98 for(unsigned I = 0; I < N; ++I);99 // CHECK: acc.serial combined(loop) {100 // CHECK: acc.loop combined(serial) {101 // CHECK: acc.yield102 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<nvidia>, #acc.device_type<radeon>], seq = [#acc.device_type<none>]} loc103 // CHECK: acc.yield104 // CHECK-NEXT: } loc105#pragma acc kernels loop auto device_type(nvidia, radeon)106 for(unsigned I = 0; I < N; ++I);107 // CHECK: acc.kernels combined(loop) {108 // CHECK: acc.loop combined(kernels) {109 // CHECK: acc.yield110 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc111 // CHECK: acc.terminator112 // CHECK-NEXT: } loc113 114#pragma acc parallel loop independent115 for(unsigned I = 0; I < N; ++I);116 // CHECK: acc.parallel combined(loop) {117 // CHECK: acc.loop combined(parallel) {118 // CHECK: acc.yield119 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc120 // CHECK: acc.yield121 // CHECK-NEXT: } loc122#pragma acc serial loop device_type(nvidia, radeon) independent123 for(unsigned I = 0; I < N; ++I);124 // CHECK: acc.serial combined(loop) {125 // CHECK: acc.loop combined(serial) {126 // CHECK: acc.yield127 // CHECK-NEXT: } attributes {independent = [#acc.device_type<nvidia>, #acc.device_type<radeon>], seq = [#acc.device_type<none>]} loc128 // CHECK: acc.yield129 // CHECK-NEXT: } loc130#pragma acc kernels loop independent device_type(nvidia, radeon)131 for(unsigned I = 0; I < N; ++I);132 // CHECK: acc.kernels combined(loop) {133 // CHECK: acc.loop combined(kernels) {134 // CHECK: acc.yield135 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc136 // CHECK: acc.terminator137 // CHECK-NEXT: } loc138 139 #pragma acc parallel loop collapse(1) device_type(radeon)140 for(unsigned I = 0; I < N; ++I)141 for(unsigned J = 0; J < N; ++J)142 for(unsigned K = 0; K < N; ++K);143 // CHECK: acc.parallel combined(loop) {144 // CHECK: acc.loop combined(parallel) {145 // CHECK: acc.yield146 // CHECK-NEXT: } attributes {collapse = [1], collapseDeviceType = [#acc.device_type<none>], independent = [#acc.device_type<none>]}147 // CHECK: acc.yield148 // CHECK-NEXT: } loc149 150 #pragma acc serial loop collapse(1) device_type(radeon) collapse (2)151 for(unsigned I = 0; I < N; ++I)152 for(unsigned J = 0; J < N; ++J)153 for(unsigned K = 0; K < N; ++K);154 // CHECK: acc.serial combined(loop) {155 // CHECK: acc.loop combined(serial) {156 // CHECK: acc.yield157 // CHECK-NEXT: } attributes {collapse = [1, 2], collapseDeviceType = [#acc.device_type<none>, #acc.device_type<radeon>], seq = [#acc.device_type<none>]}158 // CHECK: acc.yield159 // CHECK-NEXT: } loc160 161 #pragma acc kernels loop collapse(1) device_type(radeon, nvidia) collapse (2)162 for(unsigned I = 0; I < N; ++I)163 for(unsigned J = 0; J < N; ++J)164 for(unsigned K = 0; K < N; ++K);165 // CHECK: acc.kernels combined(loop) {166 // CHECK: acc.loop combined(kernels) {167 // CHECK: acc.yield168 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>], collapse = [1, 2, 2], collapseDeviceType = [#acc.device_type<none>, #acc.device_type<radeon>, #acc.device_type<nvidia>]}169 // CHECK: acc.terminator170 // CHECK-NEXT: } loc171 #pragma acc parallel loop collapse(1) device_type(radeon, nvidia) collapse(2) device_type(host) collapse(3)172 for(unsigned I = 0; I < N; ++I)173 for(unsigned J = 0; J < N; ++J)174 for(unsigned K = 0; K < N; ++K);175 // CHECK: acc.parallel combined(loop) {176 // CHECK: acc.loop combined(parallel) {177 // CHECK: acc.yield178 // CHECK-NEXT: } attributes {collapse = [1, 2, 2, 3], collapseDeviceType = [#acc.device_type<none>, #acc.device_type<radeon>, #acc.device_type<nvidia>, #acc.device_type<host>], independent = [#acc.device_type<none>]}179 // CHECK: acc.yield180 // CHECK-NEXT: } loc181 182#pragma acc kernels loop self183 for(unsigned I = 0; I < N; ++I);184 // CHECK-NEXT: acc.kernels combined(loop) {185 // CHECK-NEXT: acc.loop combined(kernels) {186 // CHECK: acc.yield187 // CHECK-NEXT: } loc188 // CHECK-NEXT: acc.terminator189 // CHECK-NEXT: } attributes {selfAttr}190 191#pragma acc serial loop self(N)192 for(unsigned I = 0; I < N; ++I);193 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i194 // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast int_to_bool %[[N_LOAD]] : !s32i -> !cir.bool195 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1196 // CHECK-NEXT: acc.serial combined(loop) self(%[[CONV_CAST]]) {197 // CHECK-NEXT: acc.loop combined(serial) {198 // CHECK: acc.yield199 // CHECK-NEXT: } loc200 // CHECK-NEXT: acc.yield201 // CHECK-NEXT: } loc202 203#pragma acc parallel loop if(N)204 for(unsigned I = 0; I < N; ++I);205 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i206 // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast int_to_bool %[[N_LOAD]] : !s32i -> !cir.bool207 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1208 // CHECK-NEXT: acc.parallel combined(loop) if(%[[CONV_CAST]]) {209 // CHECK-NEXT: acc.loop combined(parallel) {210 // CHECK: acc.yield211 // CHECK-NEXT: } loc212 // CHECK-NEXT: acc.yield213 // CHECK-NEXT: } loc214 215#pragma acc serial loop if(1)216 for(unsigned I = 0; I < N; ++I);217 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i218 // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast int_to_bool %[[ONE_LITERAL]] : !s32i -> !cir.bool219 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1220 // CHECK-NEXT: acc.serial combined(loop) if(%[[CONV_CAST]]) {221 // CHECK-NEXT: acc.loop combined(serial) {222 // CHECK: acc.yield223 // CHECK-NEXT: } loc224 // CHECK-NEXT: acc.yield225 // CHECK-NEXT: } loc226 227#pragma acc kernels loop if(N == 1)228 for(unsigned I = 0; I < N; ++I);229 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i230 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i231 // CHECK-NEXT: %[[EQ_RES:.*]] = cir.cmp(eq, %[[N_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool232 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES]] : !cir.bool to i1233 // CHECK-NEXT: acc.kernels combined(loop) if(%[[CONV_CAST]]) {234 // CHECK-NEXT: acc.loop combined(kernels) {235 // CHECK: acc.yield236 // CHECK-NEXT: } loc237 // CHECK-NEXT: acc.terminator238 // CHECK-NEXT: } loc239 240#pragma acc parallel loop if(N == 1) self(N == 2)241 for(unsigned I = 0; I < N; ++I);242 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i243 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i244 // CHECK-NEXT: %[[EQ_RES_IF:.*]] = cir.cmp(eq, %[[N_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool245 // CHECK-NEXT: %[[CONV_CAST_IF:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES_IF]] : !cir.bool to i1246 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i247 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i248 // CHECK-NEXT: %[[EQ_RES_SELF:.*]] = cir.cmp(eq, %[[N_LOAD]], %[[TWO_LITERAL]]) : !s32i, !cir.bool249 // CHECK-NEXT: %[[CONV_CAST_SELF:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES_SELF]] : !cir.bool to i1250 // CHECK-NEXT: acc.parallel combined(loop) self(%[[CONV_CAST_SELF]]) if(%[[CONV_CAST_IF]]) {251 // CHECK-NEXT: acc.loop combined(parallel) {252 // CHECK: acc.yield253 // CHECK-NEXT: } loc254 // CHECK-NEXT: acc.yield255 // CHECK-NEXT: } loc256 257 #pragma acc parallel loop tile(1, 2, 3)258 for(unsigned I = 0; I < N; ++I)259 for(unsigned J = 0; J < N; ++J)260 for(unsigned K = 0; K < N; ++K);261 // CHECK-NEXT: acc.parallel combined(loop) {262 // CHECK: %[[ONE_CONST:.*]] = arith.constant 1 : i64263 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64264 // CHECK-NEXT: %[[THREE_CONST:.*]] = arith.constant 3 : i64265 // CHECK-NEXT: acc.loop combined(parallel) tile({%[[ONE_CONST]] : i64, %[[TWO_CONST]] : i64, %[[THREE_CONST]] : i64}) {266 // CHECK: acc.yield267 // CHECK-NEXT: } loc268 // CHECK-NEXT: acc.yield269 // CHECK-NEXT: } loc270 #pragma acc serial loop tile(2) device_type(radeon)271 for(unsigned I = 0; I < N; ++I)272 for(unsigned J = 0; J < N; ++J)273 for(unsigned K = 0; K < N; ++K);274 // CHECK-NEXT: acc.serial combined(loop) {275 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64276 // CHECK-NEXT: acc.loop combined(serial) tile({%[[TWO_CONST]] : i64}) {277 // CHECK: acc.yield278 // CHECK-NEXT: } loc279 // CHECK-NEXT: acc.yield280 // CHECK-NEXT: } loc281 #pragma acc kernels loop tile(2) device_type(radeon) tile (1, *)282 for(unsigned I = 0; I < N; ++I)283 for(unsigned J = 0; J < N; ++J)284 for(unsigned K = 0; K < N; ++K);285 // CHECK-NEXT: acc.kernels combined(loop) {286 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64287 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64288 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64289 // CHECK-NEXT: acc.loop combined(kernels) tile({%[[TWO_CONST]] : i64}, {%[[ONE_CONST]] : i64, %[[STAR_CONST]] : i64} [#acc.device_type<radeon>]) {290 // CHECK: acc.yield291 // CHECK-NEXT: } loc292 // CHECK-NEXT: acc.terminator293 // CHECK-NEXT: } loc294 #pragma acc parallel loop tile(*) device_type(radeon, nvidia) tile (1, 2)295 for(unsigned I = 0; I < N; ++I)296 for(unsigned J = 0; J < N; ++J)297 for(unsigned K = 0; K < N; ++K);298 // CHECK-NEXT: acc.parallel combined(loop) {299 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64300 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64301 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64302 // CHECK-NEXT: acc.loop combined(parallel) tile({%[[STAR_CONST]] : i64}, {%[[ONE_CONST]] : i64, %[[TWO_CONST]] : i64} [#acc.device_type<radeon>], {%[[ONE_CONST]] : i64, %[[TWO_CONST]] : i64} [#acc.device_type<nvidia>]) {303 // CHECK: acc.yield304 // CHECK-NEXT: } loc305 // CHECK-NEXT: acc.yield306 // CHECK-NEXT: } loc307 #pragma acc serial loop tile(1) device_type(radeon, nvidia) tile(2, 3) device_type(host) tile(*, *, *)308 for(unsigned I = 0; I < N; ++I)309 for(unsigned J = 0; J < N; ++J)310 for(unsigned K = 0; K < N; ++K);311 // CHECK-NEXT: acc.serial combined(loop) {312 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64313 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64314 // CHECK-NEXT: %[[THREE_CONST:.*]] = arith.constant 3 : i64315 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64316 // CHECK-NEXT: %[[STAR2_CONST:.*]] = arith.constant -1 : i64317 // CHECK-NEXT: %[[STAR3_CONST:.*]] = arith.constant -1 : i64318 // CHECK-NEXT: acc.loop combined(serial) tile({%[[ONE_CONST]] : i64}, {%[[TWO_CONST]] : i64, %[[THREE_CONST]] : i64} [#acc.device_type<radeon>], {%[[TWO_CONST]] : i64, %[[THREE_CONST]] : i64} [#acc.device_type<nvidia>], {%[[STAR_CONST]] : i64, %[[STAR2_CONST]] : i64, %[[STAR3_CONST]] : i64} [#acc.device_type<host>]) {319 // CHECK: acc.yield320 // CHECK-NEXT: } loc321 // CHECK-NEXT: acc.yield322 // CHECK-NEXT: } loc323 324#pragma acc parallel loop gang325 for(unsigned I = 0; I < N; ++I);326 // CHECK-NEXT: acc.parallel combined(loop) {327 // CHECK-NEXT: acc.loop combined(parallel) gang {328 // CHECK: acc.yield329 // CHECK-NEXT: } loc330 // CHECK-NEXT: acc.yield331 // CHECK-NEXT: } loc332#pragma acc parallel loop gang device_type(nvidia) gang333 for(unsigned I = 0; I < N; ++I);334 // CHECK-NEXT: acc.parallel combined(loop) {335 // CHECK-NEXT: acc.loop combined(parallel) gang([#acc.device_type<none>, #acc.device_type<nvidia>]) {336 // CHECK: acc.yield337 // CHECK-NEXT: } loc338 // CHECK-NEXT: acc.yield339 // CHECK-NEXT: } loc340#pragma acc parallel loop gang(dim:1) device_type(nvidia) gang(dim:2)341 for(unsigned I = 0; I < N; ++I);342 // CHECK-NEXT: acc.parallel combined(loop) {343 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64344 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64345 // CHECK-NEXT: acc.loop combined(parallel) gang({dim=%[[ONE_CONST]] : i64}, {dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>]) {346 // CHECK: acc.yield347 // CHECK-NEXT: } loc348 // CHECK-NEXT: acc.yield349 // CHECK-NEXT: } loc350#pragma acc parallel loop gang(static:N, dim: 1) device_type(nvidia, radeon) gang(static:*, dim : 2)351 for(unsigned I = 0; I < N; ++I);352 // CHECK-NEXT: acc.parallel combined(loop) {353 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i354 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32355 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64356 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64357 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64358 // CHECK-NEXT: acc.loop combined(parallel) gang({static=%[[N_CONV]] : si32, dim=%[[ONE_CONST]] : i64}, {static=%[[STAR_CONST]] : i64, dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>], {static=%[[STAR_CONST]] : i64, dim=%[[TWO_CONST]] : i64} [#acc.device_type<radeon>]) {359 // CHECK: acc.yield360 // CHECK-NEXT: } loc361 // CHECK-NEXT: acc.yield362 // CHECK-NEXT: } loc363 364#pragma acc kernels loop gang(num:N) device_type(nvidia, radeon) gang(num:N)365 for(unsigned I = 0; I < N; ++I);366 // CHECK-NEXT: acc.kernels combined(loop) {367 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i368 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32369 // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i370 // CHECK-NEXT: %[[N_CONV2:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD2]] : !s32i to si32371 // CHECK-NEXT: acc.loop combined(kernels) gang({num=%[[N_CONV]] : si32}, {num=%[[N_CONV2]] : si32} [#acc.device_type<nvidia>], {num=%[[N_CONV2]] : si32} [#acc.device_type<radeon>]) {372 // CHECK: acc.yield373 // CHECK-NEXT: } loc374 // CHECK-NEXT: acc.terminator375 // CHECK-NEXT: } loc376#pragma acc kernels loop gang(static:N) device_type(nvidia) gang(static:*)377 for(unsigned I = 0; I < N; ++I);378 // CHECK-NEXT: acc.kernels combined(loop) {379 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i380 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32381 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64382 // CHECK-NEXT: acc.loop combined(kernels) gang({static=%[[N_CONV]] : si32}, {static=%[[STAR_CONST]] : i64} [#acc.device_type<nvidia>]) {383 // CHECK: acc.yield384 // CHECK-NEXT: } loc385 // CHECK-NEXT: acc.terminator386 // CHECK-NEXT: } loc387#pragma acc kernels loop gang(static:N, num: N + 1) device_type(nvidia) gang(static:*, num : N + 2)388 for(unsigned I = 0; I < N; ++I);389 // CHECK-NEXT: acc.kernels combined(loop) {390 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i391 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32392 // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i393 // CHECK-NEXT: %[[CIR_ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i394 // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD2]], %[[CIR_ONE_CONST]]) nsw : !s32i395 // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32396 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64397 // CHECK-NEXT: %[[N_LOAD3:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i398 // CHECK-NEXT: %[[CIR_TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i399 // CHECK-NEXT: %[[N_PLUS_TWO:.*]] = cir.binop(add, %[[N_LOAD3]], %[[CIR_TWO_CONST]]) nsw : !s32i400 // CHECK-NEXT: %[[N_PLUS_TWO_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_TWO]] : !s32i to si32401 // CHECK-NEXT: acc.loop combined(kernels) gang({static=%[[N_CONV]] : si32, num=%[[N_PLUS_ONE_CONV]] : si32}, {static=%[[STAR_CONST]] : i64, num=%[[N_PLUS_TWO_CONV]] : si32} [#acc.device_type<nvidia>]) {402 // CHECK: acc.yield403 // CHECK-NEXT: } loc404 // CHECK-NEXT: acc.terminator405 // CHECK-NEXT: } loc406 407#pragma acc kernels loop worker408 for(unsigned I = 0; I < N; ++I);409 // CHECK-NEXT: acc.kernels combined(loop) {410 // CHECK-NEXT: acc.loop combined(kernels) worker {411 // CHECK: acc.yield412 // CHECK-NEXT: } loc413 // CHECK: acc.terminator414 // CHECK-NEXT: } loc415 416#pragma acc kernels loop worker(N)417 for(unsigned I = 0; I < N; ++I);418 // CHECK-NEXT: acc.kernels combined(loop) {419 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i420 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32421 // CHECK-NEXT: acc.loop combined(kernels) worker(%[[N_CONV]] : si32) {422 // CHECK: acc.yield423 // CHECK-NEXT: } loc424 // CHECK: acc.terminator425 // CHECK-NEXT: } loc426 427#pragma acc kernels loop worker device_type(nvidia, radeon) worker428 for(unsigned I = 0; I < N; ++I);429 // CHECK-NEXT: acc.kernels combined(loop) {430 // CHECK-NEXT: acc.loop combined(kernels) worker([#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>]) {431 // CHECK: acc.yield432 // CHECK-NEXT: } loc433 // CHECK: acc.terminator434 // CHECK-NEXT: } loc435 436#pragma acc kernels loop worker(N) device_type(nvidia, radeon) worker437 for(unsigned I = 0; I < N; ++I);438 // CHECK-NEXT: acc.kernels combined(loop) {439 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i440 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32441 // CHECK-NEXT: acc.loop combined(kernels) worker([#acc.device_type<nvidia>, #acc.device_type<radeon>], %[[N_CONV]] : si32) {442 // CHECK: acc.yield443 // CHECK-NEXT: } loc444 // CHECK: acc.terminator445 // CHECK-NEXT: } loc446 447#pragma acc kernels loop worker device_type(nvidia, radeon) worker(N)448 for(unsigned I = 0; I < N; ++I);449 // CHECK-NEXT: acc.kernels combined(loop) {450 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i451 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32452 // CHECK-NEXT: acc.loop combined(kernels) worker([#acc.device_type<none>], %[[N_CONV]] : si32 [#acc.device_type<nvidia>], %[[N_CONV]] : si32 [#acc.device_type<radeon>]) {453 // CHECK: acc.yield454 // CHECK-NEXT: } loc455 // CHECK: acc.terminator456 // CHECK-NEXT: } loc457 458#pragma acc kernels loop worker(N) device_type(nvidia, radeon) worker(N + 1)459 for(unsigned I = 0; I < N; ++I);460 // CHECK-NEXT: acc.kernels combined(loop) {461 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i462 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32463 // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i464 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i465 // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD2]], %[[ONE_CONST]]) nsw : !s32i466 // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32467 // CHECK-NEXT: acc.loop combined(kernels) worker(%[[N_CONV]] : si32, %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<nvidia>], %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<radeon>]) {468 // CHECK: acc.yield469 // CHECK-NEXT: } loc470 // CHECK: acc.terminator471 // CHECK-NEXT: } loc472 473#pragma acc kernels loop device_type(nvidia, radeon) worker(num:N + 1)474 for(unsigned I = 0; I < N; ++I);475 // CHECK-NEXT: acc.kernels combined(loop) {476 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i477 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i478 // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD]], %[[ONE_CONST]]) nsw : !s32i479 // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32480 // CHECK-NEXT: acc.loop combined(kernels) worker(%[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<nvidia>], %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<radeon>]) {481 // CHECK: acc.terminator482 // CHECK-NEXT: } loc483 484 485#pragma acc kernels loop worker vector device_type(nvidia) worker vector486 for(unsigned I = 0; I < N; ++I);487 // CHECK-NEXT: acc.kernels combined(loop) {488 // CHECK-NEXT: acc.loop combined(kernels) worker([#acc.device_type<none>, #acc.device_type<nvidia>]) vector([#acc.device_type<none>, #acc.device_type<nvidia>])489 // CHECK: acc.yield490 // CHECK-NEXT: } loc491 // CHECK: acc.terminator492 // CHECK-NEXT: } loc493 494#pragma acc kernels loop vector495 for(unsigned I = 0; I < N; ++I);496 // CHECK-NEXT: acc.kernels combined(loop) {497 // CHECK: acc.loop combined(kernels) vector {498 // CHECK: acc.yield499 // CHECK-NEXT: } loc500 // CHECK: acc.terminator501 // CHECK-NEXT: } loc502 503#pragma acc kernels loop vector(N)504 for(unsigned I = 0; I < N; ++I);505 // CHECK-NEXT: acc.kernels combined(loop) {506 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i507 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32508 // CHECK-NEXT: acc.loop combined(kernels) vector(%[[N_CONV]] : si32) {509 // CHECK: acc.yield510 // CHECK-NEXT: } loc511 // CHECK: acc.terminator512 // CHECK-NEXT: } loc513 514#pragma acc kernels loop vector device_type(nvidia, radeon) vector515 for(unsigned I = 0; I < N; ++I);516 // CHECK-NEXT: acc.kernels combined(loop) {517 // CHECK-NEXT: acc.loop combined(kernels) vector([#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>]) {518 // CHECK: acc.yield519 // CHECK-NEXT: } loc520 // CHECK: acc.terminator521 // CHECK-NEXT: } loc522 523#pragma acc kernels loop vector(N) device_type(nvidia, radeon) vector524 for(unsigned I = 0; I < N; ++I);525 // CHECK-NEXT: acc.kernels combined(loop) {526 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i527 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32528 // CHECK-NEXT: acc.loop combined(kernels) vector([#acc.device_type<nvidia>, #acc.device_type<radeon>], %[[N_CONV]] : si32) {529 // CHECK: acc.yield530 // CHECK-NEXT: } loc531 // CHECK: acc.terminator532 // CHECK-NEXT: } loc533 534#pragma acc kernels loop vector(N) device_type(nvidia, radeon) vector(N + 1)535 for(unsigned I = 0; I < N; ++I);536 // CHECK-NEXT: acc.kernels combined(loop) {537 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i538 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32539 // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i540 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i541 // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD2]], %[[ONE_CONST]]) nsw : !s32i542 // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32543 // CHECK-NEXT: acc.loop combined(kernels) vector(%[[N_CONV]] : si32, %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<nvidia>], %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<radeon>]) {544 // CHECK: acc.yield545 // CHECK-NEXT: } loc546 // CHECK: acc.terminator547 // CHECK-NEXT: } loc548 549#pragma acc kernels loop device_type(nvidia, radeon) vector(length:N + 1)550 for(unsigned I = 0; I < N; ++I);551 // CHECK-NEXT: acc.kernels combined(loop) {552 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i553 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i554 // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD]], %[[ONE_CONST]]) nsw : !s32i555 // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32556 // CHECK-NEXT: acc.loop combined(kernels) vector(%[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<nvidia>], %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<radeon>]) {557 // CHECK: acc.yield558 // CHECK-NEXT: } loc559 // CHECK: acc.terminator560 // CHECK-NEXT: } loc561 562#pragma acc kernels loop worker(N) vector(N) device_type(nvidia) worker(N) vector(N)563 for(unsigned I = 0; I < N; ++I);564 // CHECK-NEXT: acc.kernels combined(loop) {565 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i566 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32567 // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i568 // CHECK-NEXT: %[[N_CONV2:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD2]] : !s32i to si32569 // CHECK-NEXT: %[[N_LOAD3:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i570 // CHECK-NEXT: %[[N_CONV3:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD3]] : !s32i to si32571 // CHECK-NEXT: %[[N_LOAD4:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i572 // CHECK-NEXT: %[[N_CONV4:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD4]] : !s32i to si32573 // CHECK-NEXT: acc.loop combined(kernels) worker(%[[N_CONV]] : si32, %[[N_CONV3]] : si32 [#acc.device_type<nvidia>]) vector(%[[N_CONV2]] : si32, %[[N_CONV4]] : si32 [#acc.device_type<nvidia>]) {574 // CHECK: acc.yield575 // CHECK-NEXT: } loc576 // CHECK: acc.terminator577 // CHECK-NEXT: } loc578 579#pragma acc parallel loop wait580 for(unsigned I = 0; I < N; ++I);581 // CHECK-NEXT: acc.parallel combined(loop) wait {582 // CHECK-NEXT: acc.loop combined(parallel) {583 // CHECK: acc.yield584 // CHECK-NEXT: } loc585 // CHECK-NEXT: acc.yield586 // CHECK-NEXT: } loc587 588#pragma acc serial loop wait device_type(nvidia) wait589 for(unsigned I = 0; I < N; ++I);590 // CHECK-NEXT: acc.serial combined(loop) wait([#acc.device_type<none>, #acc.device_type<nvidia>]) {591 // CHECK-NEXT: acc.loop combined(serial) {592 // CHECK: acc.yield593 // CHECK-NEXT: } loc594 // CHECK-NEXT: acc.yield595 // CHECK-NEXT: } loc596 597#pragma acc kernels loop wait(1) device_type(nvidia) wait598 for(unsigned I = 0; I < N; ++I);599 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i600 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32601 // CHECK-NEXT: acc.kernels combined(loop) wait([#acc.device_type<nvidia>], {%[[ONE_CAST]] : si32}) {602 // CHECK-NEXT: acc.loop combined(kernels) {603 // CHECK: acc.yield604 // CHECK-NEXT: } loc605 // CHECK-NEXT: acc.terminator606 // CHECK-NEXT: } loc607 608#pragma acc parallel loop wait device_type(nvidia) wait(1)609 for(unsigned I = 0; I < N; ++I);610 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i611 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32612 // CHECK-NEXT: acc.parallel combined(loop) wait([#acc.device_type<none>], {%[[ONE_CAST]] : si32} [#acc.device_type<nvidia>]) {613 // CHECK-NEXT: acc.loop combined(parallel) {614 // CHECK: acc.yield615 // CHECK-NEXT: } loc616 // CHECK-NEXT: acc.yield617 // CHECK-NEXT: } loc618 619#pragma acc serial loop wait(1) device_type(nvidia) wait(1)620 for(unsigned I = 0; I < N; ++I);621 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i622 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32623 // CHECK-NEXT: %[[ONE_LITERAL2:.*]] = cir.const #cir.int<1> : !s32i624 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL2]] : !s32i to si32625 // CHECK-NEXT: acc.serial combined(loop) wait({%[[ONE_CAST]] : si32}, {%[[ONE_CAST2]] : si32} [#acc.device_type<nvidia>]) {626 // CHECK-NEXT: acc.loop combined(serial) {627 // CHECK: acc.yield628 // CHECK-NEXT: } loc629 // CHECK-NEXT: acc.yield630 // CHECK-NEXT: } loc631 632#pragma acc kernels loop wait(devnum: cond : 1)633 for(unsigned I = 0; I < N; ++I);634 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i635 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32636 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i637 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32638 // CHECK-NEXT: acc.kernels combined(loop) wait({devnum: %[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32}) {639 // CHECK-NEXT: acc.loop combined(kernels) {640 // CHECK: acc.yield641 // CHECK-NEXT: } loc642 // CHECK-NEXT: acc.terminator643 // CHECK-NEXT: } loc644 645#pragma acc parallel loop wait(devnum: cond : 1) device_type(nvidia) wait(devnum: cond : 1)646 for(unsigned I = 0; I < N; ++I);647 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i648 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32649 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i650 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32651 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i652 // CHECK-NEXT: %[[CONV_CAST2:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32653 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i654 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32655 // CHECK-NEXT: acc.parallel combined(loop) wait({devnum: %[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32}, {devnum: %[[CONV_CAST2]] : si32, %[[ONE_CAST2]] : si32} [#acc.device_type<nvidia>]) {656 // CHECK-NEXT: acc.loop combined(parallel) {657 // CHECK: acc.yield658 // CHECK-NEXT: } loc659 // CHECK-NEXT: acc.yield660 // CHECK-NEXT: } loc661 662#pragma acc serial loop wait(devnum: cond : 1, 2)663 for(unsigned I = 0; I < N; ++I);664 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i665 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32666 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i667 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32668 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i669 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32670 // CHECK-NEXT: acc.serial combined(loop) wait({devnum: %[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32, %[[TWO_CAST]] : si32}) {671 // CHECK-NEXT: acc.loop combined(serial) {672 // CHECK: acc.yield673 // CHECK-NEXT: } loc674 // CHECK-NEXT: acc.yield675 // CHECK-NEXT: } loc676 677#pragma acc kernels loop wait(devnum: cond : 1, 2) device_type(nvidia, radeon) wait(devnum: cond : 1, 2)678 for(unsigned I = 0; I < N; ++I);679 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i680 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32681 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i682 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32683 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i684 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32685 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i686 // CHECK-NEXT: %[[CONV_CAST2:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32687 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i688 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32689 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i690 // CHECK-NEXT: %[[TWO_CAST2:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32691 // CHECK-NEXT: acc.kernels combined(loop) wait({devnum: %[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32, %[[TWO_CAST]] : si32}, {devnum: %[[CONV_CAST2]] : si32, %[[ONE_CAST2]] : si32, %[[TWO_CAST2]] : si32} [#acc.device_type<nvidia>], {devnum: %[[CONV_CAST2]] : si32, %[[ONE_CAST2]] : si32, %[[TWO_CAST2]] : si32} [#acc.device_type<radeon>]) {692 // CHECK-NEXT: acc.loop combined(kernels) {693 // CHECK: acc.yield694 // CHECK-NEXT: } loc695 // CHECK-NEXT: acc.terminator696 // CHECK-NEXT: } loc697 698#pragma acc parallel loop wait(cond, 1)699 for(unsigned I = 0; I < N; ++I);700 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i701 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32702 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i703 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32704 // CHECK-NEXT: acc.parallel combined(loop) wait({%[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32}) {705 // CHECK-NEXT: acc.loop combined(parallel) {706 // CHECK: acc.yield707 // CHECK-NEXT: } loc708 // CHECK-NEXT: acc.yield709 // CHECK-NEXT: } loc710 711#pragma acc serial loop wait(queues: cond, 1) device_type(radeon)712 for(unsigned I = 0; I < N; ++I);713 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i714 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32715 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i716 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32717 // CHECK-NEXT: acc.serial combined(loop) wait({%[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32}) {718 // CHECK-NEXT: acc.loop combined(serial) {719 // CHECK: acc.yield720 // CHECK-NEXT: } loc721 // CHECK-NEXT: acc.yield722 // CHECK-NEXT: } loc723 724#pragma acc parallel loop num_gangs(1)725 for(unsigned I = 0; I < N; ++I);726 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i727 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32728 // CHECK-NEXT: acc.parallel combined(loop) num_gangs({%[[ONE_CAST]] : si32}) {729 // CHECK-NEXT: acc.loop combined(parallel) {730 // CHECK: acc.yield731 // CHECK-NEXT: } loc732 // CHECK-NEXT: acc.yield733 // CHECK-NEXT: } loc734 735#pragma acc kernels loop num_gangs(cond)736 for(unsigned I = 0; I < N; ++I);737 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i738 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32739 // CHECK-NEXT: acc.kernels combined(loop) num_gangs({%[[CONV_CAST]] : si32}) {740 // CHECK-NEXT: acc.loop combined(kernels) {741 // CHECK: acc.yield742 // CHECK-NEXT: } loc743 // CHECK-NEXT: acc.terminator744 // CHECK-NEXT: } loc745 746#pragma acc parallel loop num_gangs(1, cond, 2)747 for(unsigned I = 0; I < N; ++I);748 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i749 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32750 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i751 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32752 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i753 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32754 // CHECK-NEXT: acc.parallel combined(loop) num_gangs({%[[ONE_CAST]] : si32, %[[CONV_CAST]] : si32, %[[TWO_CAST]] : si32}) {755 // CHECK-NEXT: acc.loop combined(parallel) {756 // CHECK: acc.yield757 // CHECK-NEXT: } loc758 // CHECK-NEXT: acc.yield759 // CHECK-NEXT: } loc760 761#pragma acc kernels loop num_gangs(1) device_type(radeon) num_gangs(cond)762 for(unsigned I = 0; I < N; ++I);763 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i764 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32765 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i766 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32767 // CHECK-NEXT: acc.kernels combined(loop) num_gangs({%[[ONE_CAST]] : si32}, {%[[CONV_CAST]] : si32} [#acc.device_type<radeon>]) {768 // CHECK-NEXT: acc.loop combined(kernels) {769 // CHECK: acc.yield770 // CHECK-NEXT: } loc771 // CHECK-NEXT: acc.terminator772 // CHECK-NEXT: } loc773 774#pragma acc parallel loop num_gangs(1, cond, 2) device_type(radeon) num_gangs(4, 5, 6)775 for(unsigned I = 0; I < N; ++I);776 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i777 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32778 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i779 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32780 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i781 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32782 // CHECK-NEXT: %[[FOUR_LITERAL:.*]] = cir.const #cir.int<4> : !s32i783 // CHECK-NEXT: %[[FOUR_CAST:.*]] = builtin.unrealized_conversion_cast %[[FOUR_LITERAL]] : !s32i to si32784 // CHECK-NEXT: %[[FIVE_LITERAL:.*]] = cir.const #cir.int<5> : !s32i785 // CHECK-NEXT: %[[FIVE_CAST:.*]] = builtin.unrealized_conversion_cast %[[FIVE_LITERAL]] : !s32i to si32786 // CHECK-NEXT: %[[SIX_LITERAL:.*]] = cir.const #cir.int<6> : !s32i787 // CHECK-NEXT: %[[SIX_CAST:.*]] = builtin.unrealized_conversion_cast %[[SIX_LITERAL]] : !s32i to si32788 // CHECK-NEXT: acc.parallel combined(loop) num_gangs({%[[ONE_CAST]] : si32, %[[CONV_CAST]] : si32, %[[TWO_CAST]] : si32}, {%[[FOUR_CAST]] : si32, %[[FIVE_CAST]] : si32, %[[SIX_CAST]] : si32} [#acc.device_type<radeon>])789 // CHECK-NEXT: acc.loop combined(parallel) {790 // CHECK: acc.yield791 // CHECK-NEXT: } loc792 // CHECK-NEXT: acc.yield793 // CHECK-NEXT: } loc794 795#pragma acc parallel loop num_gangs(1, cond, 2) device_type(radeon, nvidia) num_gangs(4, 5, 6)796 for(unsigned I = 0; I < N; ++I);797 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i798 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32799 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i800 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32801 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i802 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32803 // CHECK-NEXT: %[[FOUR_LITERAL:.*]] = cir.const #cir.int<4> : !s32i804 // CHECK-NEXT: %[[FOUR_CAST:.*]] = builtin.unrealized_conversion_cast %[[FOUR_LITERAL]] : !s32i to si32805 // CHECK-NEXT: %[[FIVE_LITERAL:.*]] = cir.const #cir.int<5> : !s32i806 // CHECK-NEXT: %[[FIVE_CAST:.*]] = builtin.unrealized_conversion_cast %[[FIVE_LITERAL]] : !s32i to si32807 // CHECK-NEXT: %[[SIX_LITERAL:.*]] = cir.const #cir.int<6> : !s32i808 // CHECK-NEXT: %[[SIX_CAST:.*]] = builtin.unrealized_conversion_cast %[[SIX_LITERAL]] : !s32i to si32809 // CHECK-NEXT: acc.parallel combined(loop) num_gangs({%[[ONE_CAST]] : si32, %[[CONV_CAST]] : si32, %[[TWO_CAST]] : si32}, {%[[FOUR_CAST]] : si32, %[[FIVE_CAST]] : si32, %[[SIX_CAST]] : si32} [#acc.device_type<radeon>], {%[[FOUR_CAST]] : si32, %[[FIVE_CAST]] : si32, %[[SIX_CAST]] : si32} [#acc.device_type<nvidia>])810 // CHECK-NEXT: acc.loop combined(parallel) {811 // CHECK: acc.yield812 // CHECK-NEXT: } loc813 // CHECK-NEXT: acc.yield814 // CHECK-NEXT: } loc815 816#pragma acc parallel loop num_workers(cond)817 for(unsigned I = 0; I < N; ++I);818 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i819 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32820 // CHECK-NEXT: acc.parallel combined(loop) num_workers(%[[CONV_CAST]] : si32) {821 // CHECK-NEXT: acc.loop combined(parallel) {822 // CHECK: acc.yield823 // CHECK-NEXT: } loc824 // CHECK-NEXT: acc.yield825 // CHECK-NEXT: } loc826 827#pragma acc kernels loop num_workers(cond) device_type(nvidia) num_workers(2u)828 for(unsigned I = 0; I < N; ++I);829 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i830 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32831 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !u32i832 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !u32i to ui32833 // CHECK-NEXT: acc.kernels combined(loop) num_workers(%[[CONV_CAST]] : si32, %[[TWO_CAST]] : ui32 [#acc.device_type<nvidia>]) {834 // CHECK-NEXT: acc.loop combined(kernels) {835 // CHECK: acc.yield836 // CHECK-NEXT: } loc837 // CHECK-NEXT: acc.terminator838 // CHECK-NEXT: } loc839 840#pragma acc parallel loop num_workers(cond) device_type(nvidia, host) num_workers(2) device_type(radeon) num_workers(3)841 for(unsigned I = 0; I < N; ++I);842 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i843 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32844 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i845 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32846 // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i847 // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32848 // CHECK-NEXT: acc.parallel combined(loop) num_workers(%[[CONV_CAST]] : si32, %[[TWO_CAST]] : si32 [#acc.device_type<nvidia>], %[[TWO_CAST]] : si32 [#acc.device_type<host>], %[[THREE_CAST]] : si32 [#acc.device_type<radeon>]) {849 // CHECK-NEXT: acc.loop combined(parallel) {850 // CHECK: acc.yield851 // CHECK-NEXT: } loc852 // CHECK-NEXT: acc.yield853 // CHECK-NEXT: } loc854 855#pragma acc kernels loop num_workers(cond) device_type(nvidia) num_workers(2) device_type(radeon, multicore) num_workers(4)856 for(unsigned I = 0; I < N; ++I);857 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i858 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32859 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i860 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32861 // CHECK-NEXT: %[[FOUR_LITERAL:.*]] = cir.const #cir.int<4> : !s32i862 // CHECK-NEXT: %[[FOUR_CAST:.*]] = builtin.unrealized_conversion_cast %[[FOUR_LITERAL]] : !s32i to si32863 // CHECK-NEXT: acc.kernels combined(loop) num_workers(%[[CONV_CAST]] : si32, %[[TWO_CAST]] : si32 [#acc.device_type<nvidia>], %[[FOUR_CAST]] : si32 [#acc.device_type<radeon>], %[[FOUR_CAST]] : si32 [#acc.device_type<multicore>]) {864 // CHECK-NEXT: acc.loop combined(kernels) {865 // CHECK: acc.yield866 // CHECK-NEXT: } loc867 // CHECK-NEXT: acc.terminator868 // CHECK-NEXT: } loc869 870#pragma acc parallel loop device_type(nvidia) num_workers(2) device_type(radeon) num_workers(3)871 for(unsigned I = 0; I < N; ++I);872 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i873 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32874 // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i875 // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32876 // CHECK-NEXT: acc.parallel combined(loop) num_workers(%[[TWO_CAST]] : si32 [#acc.device_type<nvidia>], %[[THREE_CAST]] : si32 [#acc.device_type<radeon>]) {877 // CHECK-NEXT: acc.loop combined(parallel) {878 // CHECK: acc.yield879 // CHECK-NEXT: } loc880 // CHECK-NEXT: acc.yield881 // CHECK-NEXT: } loc882 //883#pragma acc parallel loop vector_length(cond)884 for(unsigned I = 0; I < N; ++I);885 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i886 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32887 // CHECK-NEXT: acc.parallel combined(loop) vector_length(%[[CONV_CAST]] : si32) {888 // CHECK-NEXT: acc.loop combined(parallel) {889 // CHECK: acc.yield890 // CHECK-NEXT: } loc891 // CHECK-NEXT: acc.yield892 // CHECK-NEXT: } loc893 894#pragma acc kernels loop vector_length(cond) device_type(nvidia) vector_length(2u)895 for(unsigned I = 0; I < N; ++I);896 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i897 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32898 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !u32i899 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !u32i to ui32900 // CHECK-NEXT: acc.kernels combined(loop) vector_length(%[[CONV_CAST]] : si32, %[[TWO_CAST]] : ui32 [#acc.device_type<nvidia>]) {901 // CHECK-NEXT: acc.loop combined(kernels) {902 // CHECK: acc.yield903 // CHECK-NEXT: } loc904 // CHECK-NEXT: acc.terminator905 // CHECK-NEXT: } loc906 907#pragma acc parallel loop vector_length(cond) device_type(nvidia, host) vector_length(2) device_type(radeon) vector_length(3)908 for(unsigned I = 0; I < N; ++I);909 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i910 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32911 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i912 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32913 // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i914 // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32915 // CHECK-NEXT: acc.parallel combined(loop) vector_length(%[[CONV_CAST]] : si32, %[[TWO_CAST]] : si32 [#acc.device_type<nvidia>], %[[TWO_CAST]] : si32 [#acc.device_type<host>], %[[THREE_CAST]] : si32 [#acc.device_type<radeon>]) {916 // CHECK-NEXT: acc.loop combined(parallel) {917 // CHECK: acc.yield918 // CHECK-NEXT: } loc919 // CHECK-NEXT: acc.yield920 // CHECK-NEXT: } loc921 922#pragma acc kernels loop vector_length(cond) device_type(nvidia) vector_length(2) device_type(radeon, multicore) vector_length(4)923 for(unsigned I = 0; I < N; ++I);924 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i925 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32926 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i927 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32928 // CHECK-NEXT: %[[FOUR_LITERAL:.*]] = cir.const #cir.int<4> : !s32i929 // CHECK-NEXT: %[[FOUR_CAST:.*]] = builtin.unrealized_conversion_cast %[[FOUR_LITERAL]] : !s32i to si32930 // CHECK-NEXT: acc.kernels combined(loop) vector_length(%[[CONV_CAST]] : si32, %[[TWO_CAST]] : si32 [#acc.device_type<nvidia>], %[[FOUR_CAST]] : si32 [#acc.device_type<radeon>], %[[FOUR_CAST]] : si32 [#acc.device_type<multicore>]) {931 // CHECK-NEXT: acc.loop combined(kernels) {932 // CHECK: acc.yield933 // CHECK-NEXT: } loc934 // CHECK-NEXT: acc.terminator935 // CHECK-NEXT: } loc936 937#pragma acc parallel loop device_type(nvidia) vector_length(2) device_type(radeon) vector_length(3)938 for(unsigned I = 0; I < N; ++I);939 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i940 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32941 // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i942 // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32943 // CHECK-NEXT: acc.parallel combined(loop) vector_length(%[[TWO_CAST]] : si32 [#acc.device_type<nvidia>], %[[THREE_CAST]] : si32 [#acc.device_type<radeon>]) {944 // CHECK-NEXT: acc.loop combined(parallel) {945 // CHECK: acc.yield946 // CHECK-NEXT: } loc947 // CHECK-NEXT: acc.yield948 // CHECK-NEXT: } loc949 950#pragma acc parallel loop async951 for(unsigned I = 0; I < N; ++I);952 // CHECK-NEXT: acc.parallel combined(loop) async {953 // CHECK-NEXT: acc.loop combined(parallel) {954 // CHECK: acc.yield955 // CHECK-NEXT: } loc956 // CHECK-NEXT: acc.yield957 // CHECK-NEXT: } loc958 959#pragma acc serial loop async(cond)960 for(unsigned I = 0; I < N; ++I);961 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i962 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32963 // CHECK-NEXT: acc.serial combined(loop) async(%[[CONV_CAST]] : si32) {964 // CHECK-NEXT: acc.loop combined(serial) {965 // CHECK: acc.yield966 // CHECK-NEXT: } loc967 // CHECK-NEXT: acc.yield968 // CHECK-NEXT: } loc969 970#pragma acc kernels loop async device_type(nvidia, radeon) async971 for(unsigned I = 0; I < N; ++I);972 // CHECK-NEXT: acc.kernels combined(loop) async([#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>]) {973 // CHECK-NEXT: acc.loop combined(kernels) {974 // CHECK: acc.yield975 // CHECK-NEXT: } loc976 // CHECK-NEXT: acc.terminator977 // CHECK-NEXT: } loc978 979#pragma acc parallel loop async(3) device_type(nvidia, radeon) async(cond)980 for(unsigned I = 0; I < N; ++I);981 // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i982 // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si32983 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i984 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32985 // CHECK-NEXT: acc.parallel combined(loop) async(%[[THREE_CAST]] : si32, %[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {986 // CHECK-NEXT: acc.loop combined(parallel) {987 // CHECK: acc.yield988 // CHECK-NEXT: } loc989 // CHECK-NEXT: acc.yield990 // CHECK-NEXT: } loc991 992#pragma acc serial loop async device_type(nvidia, radeon) async(cond)993 for(unsigned I = 0; I < N; ++I);994 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i995 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32996 // CHECK-NEXT: acc.serial combined(loop) async([#acc.device_type<none>], %[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {997 // CHECK-NEXT: acc.loop combined(serial) {998 // CHECK: acc.yield999 // CHECK-NEXT: } loc1000 // CHECK-NEXT: acc.yield1001 // CHECK-NEXT: } loc1002 1003#pragma acc kernels loop async(3) device_type(nvidia, radeon) async1004 for(unsigned I = 0; I < N; ++I);1005 // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i1006 // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si321007 // CHECK-NEXT: acc.kernels combined(loop) async([#acc.device_type<nvidia>, #acc.device_type<radeon>], %[[THREE_CAST]] : si32) {1008 // CHECK-NEXT: acc.loop combined(kernels) {1009 // CHECK: acc.yield1010 // CHECK-NEXT: } loc1011 // CHECK-NEXT: acc.terminator1012 // CHECK-NEXT: } loc1013}1014extern "C" void acc_combined_data_clauses(int *arg1, int *arg2) {1015 // CHECK: cir.func{{.*}} @acc_combined_data_clauses(%[[ARG1_PARAM:.*]]: !cir.ptr<!s32i>{{.*}}, %[[ARG2_PARAM:.*]]: !cir.ptr<!s32i>{{.*}}) {1016 // CHECK-NEXT: %[[ARG1:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["arg1", init]1017 // CHECK-NEXT: %[[ARG2:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["arg2", init]1018 // CHECK-NEXT: cir.store %[[ARG1_PARAM]], %[[ARG1]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>1019 // CHECK-NEXT: cir.store %[[ARG2_PARAM]], %[[ARG2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>1020 1021#pragma acc parallel loop deviceptr(arg1)1022 for(unsigned I = 0; I < 5; ++I);1023 // CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}1024 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[DEVPTR1]] : !cir.ptr<!cir.ptr<!s32i>>) {1025 // CHECK-NEXT: acc.loop combined(parallel) {1026 // CHECK: acc.yield1027 // CHECK-NEXT: } loc1028 // CHECK-NEXT: acc.yield1029 // CHECK-NEXT: } loc1030 1031#pragma acc serial loop deviceptr(arg2)1032 for(unsigned I = 0; I < 5; ++I);1033 // CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}1034 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>) {1035 // CHECK-NEXT: acc.loop combined(serial) {1036 // CHECK: acc.yield1037 // CHECK-NEXT: } loc1038 // CHECK-NEXT: acc.yield1039 // CHECK-NEXT: } loc1040 1041#pragma acc kernels loop deviceptr(arg1, arg2)1042 for(unsigned I = 0; I < 5; ++I);1043 // CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}1044 // CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}1045 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) {1046 // CHECK-NEXT: acc.loop combined(kernels) {1047 // CHECK: acc.yield1048 // CHECK-NEXT: } loc1049 // CHECK-NEXT: acc.terminator1050 // CHECK-NEXT: } loc1051 1052#pragma acc parallel loop deviceptr(arg1) async1053 for(unsigned I = 0; I < 5; ++I);1054 // CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}1055 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[DEVPTR1]] : !cir.ptr<!cir.ptr<!s32i>>) async {1056 // CHECK-NEXT: acc.loop combined(parallel) {1057 // CHECK: acc.yield1058 // CHECK-NEXT: } loc1059 // CHECK-NEXT: acc.yield1060 // CHECK-NEXT: } loc1061 1062#pragma acc serial loop deviceptr(arg2) async device_type(nvidia)1063 for(unsigned I = 0; I < 5; ++I);1064 // CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}1065 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>) async {1066 // CHECK-NEXT: acc.loop combined(serial) {1067 // CHECK: acc.yield1068 // CHECK-NEXT: } loc1069 // CHECK-NEXT: acc.yield1070 // CHECK-NEXT: } loc1071 1072#pragma acc kernels loop deviceptr(arg1, arg2) device_type(nvidia) async1073 for(unsigned I = 0; I < 5; ++I);1074 // CHECK-NEXT: %[[DEVPTR1:.*]] = acc.deviceptr varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}1075 // CHECK-NEXT: %[[DEVPTR2:.*]] = acc.deviceptr varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}1076 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[DEVPTR1]], %[[DEVPTR2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<nvidia>]) {1077 // CHECK-NEXT: acc.loop combined(kernels) {1078 // CHECK: acc.yield1079 // CHECK-NEXT: } loc1080 // CHECK-NEXT: acc.terminator1081 // CHECK-NEXT: } loc1082 1083#pragma acc parallel loop no_create(arg1)1084 for(unsigned I = 0; I < 5; ++I);1085 // CHECK-NEXT: %[[NOCREATE1:.*]] = acc.nocreate varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}1086 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) {1087 // CHECK-NEXT: acc.loop combined(parallel) {1088 // CHECK: acc.yield1089 // CHECK-NEXT: } loc1090 // CHECK-NEXT: acc.yield1091 // CHECK-NEXT: } loc1092 // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_no_create>, name = "arg1"}1093 1094#pragma acc serial loop no_create(arg2)1095 for(unsigned I = 0; I < 5; ++I);1096 // CHECK-NEXT: %[[NOCREATE2:.*]] = acc.nocreate varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}1097 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) {1098 // CHECK-NEXT: acc.loop combined(serial) {1099 // CHECK: acc.yield1100 // CHECK-NEXT: } loc1101 // CHECK-NEXT: acc.yield1102 // CHECK-NEXT: } loc1103 // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_no_create>, name = "arg2"}1104 1105#pragma acc kernels loop no_create(arg1, arg2) device_type(host) async1106 for(unsigned I = 0; I < 5; ++I);1107 // CHECK-NEXT: %[[NOCREATE1:.*]] = acc.nocreate varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}1108 // CHECK-NEXT: %[[NOCREATE2:.*]] = acc.nocreate varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}1109 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[NOCREATE1]], %[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {1110 // CHECK-NEXT: acc.loop combined(kernels) {1111 // CHECK: acc.yield1112 // CHECK-NEXT: } loc1113 // CHECK-NEXT: acc.terminator1114 // CHECK-NEXT: } loc1115 // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {dataClause = #acc<data_clause acc_no_create>, name = "arg2"}1116 // CHECK-NEXT: acc.delete accPtr(%[[NOCREATE1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {dataClause = #acc<data_clause acc_no_create>, name = "arg1"}1117 1118#pragma acc parallel loop present(arg1)1119 for(unsigned I = 0; I < 5; ++I);1120 // CHECK-NEXT: %[[PRESENT1:.*]] = acc.present varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}1121 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[PRESENT1]] : !cir.ptr<!cir.ptr<!s32i>>) {1122 // CHECK-NEXT: acc.loop combined(parallel) {1123 // CHECK: acc.yield1124 // CHECK-NEXT: } loc1125 // CHECK-NEXT: acc.yield1126 // CHECK-NEXT: } loc1127 // CHECK-NEXT: acc.delete accPtr(%[[PRESENT1]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_present>, name = "arg1"}1128 1129#pragma acc serial loop present(arg2)1130 for(unsigned I = 0; I < 5; ++I);1131 // CHECK-NEXT: %[[PRESENT2:.*]] = acc.present varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}1132 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[PRESENT2]] : !cir.ptr<!cir.ptr<!s32i>>) {1133 // CHECK-NEXT: acc.loop combined(serial) {1134 // CHECK: acc.yield1135 // CHECK-NEXT: } loc1136 // CHECK-NEXT: acc.yield1137 // CHECK-NEXT: } loc1138 // CHECK-NEXT: acc.delete accPtr(%[[PRESENT2]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_present>, name = "arg2"}1139 1140#pragma acc kernels loop present(arg1, arg2) device_type(host) async1141 for(unsigned I = 0; I < 5; ++I);1142 // CHECK-NEXT: %[[PRESENT1:.*]] = acc.present varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}1143 // CHECK-NEXT: %[[PRESENT2:.*]] = acc.present varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}1144 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[PRESENT1]], %[[PRESENT2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {1145 // CHECK-NEXT: acc.loop combined(kernels) {1146 // CHECK: acc.yield1147 // CHECK-NEXT: } loc1148 // CHECK-NEXT: acc.terminator1149 // CHECK-NEXT: } loc1150 // CHECK-NEXT: acc.delete accPtr(%[[PRESENT2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {dataClause = #acc<data_clause acc_present>, name = "arg2"}1151 // CHECK-NEXT: acc.delete accPtr(%[[PRESENT1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {dataClause = #acc<data_clause acc_present>, name = "arg1"}1152 1153#pragma acc parallel loop attach(arg1)1154 for(unsigned I = 0; I < 5; ++I);1155 // CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}1156 // CHECK-NEXT: acc.parallel combined(loop) dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>) {1157 // CHECK-NEXT: acc.loop combined(parallel) {1158 // CHECK: acc.yield1159 // CHECK-NEXT: } loc1160 // CHECK-NEXT: acc.yield1161 // CHECK-NEXT: } loc1162 // CHECK-NEXT: acc.detach accPtr(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_attach>, name = "arg1"}1163 1164#pragma acc serial loop attach(arg2)1165 for(unsigned I = 0; I < 5; ++I);1166 // CHECK-NEXT: %[[ATTACH2:.*]] = acc.attach varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}1167 // CHECK-NEXT: acc.serial combined(loop) dataOperands(%[[ATTACH2]] : !cir.ptr<!cir.ptr<!s32i>>) {1168 // CHECK-NEXT: acc.loop combined(serial) {1169 // CHECK: acc.yield1170 // CHECK-NEXT: } loc1171 // CHECK-NEXT: acc.yield1172 // CHECK-NEXT: } loc1173 // CHECK-NEXT: acc.detach accPtr(%[[ATTACH2]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_attach>, name = "arg2"}1174 1175#pragma acc kernels loop attach(arg1, arg2) device_type(host) async1176 for(unsigned I = 0; I < 5; ++I);1177 // CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[ARG1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg1"}1178 // CHECK-NEXT: %[[ATTACH2:.*]] = acc.attach varPtr(%[[ARG2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "arg2"}1179 // CHECK-NEXT: acc.kernels combined(loop) dataOperands(%[[ATTACH1]], %[[ATTACH2]] : !cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {1180 // CHECK-NEXT: acc.loop combined(kernels) {1181 // CHECK: acc.yield1182 // CHECK-NEXT: } loc1183 // CHECK-NEXT: acc.terminator1184 // CHECK-NEXT: } loc1185 // CHECK-NEXT: acc.detach accPtr(%[[ATTACH2]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {dataClause = #acc<data_clause acc_attach>, name = "arg2"}1186 // CHECK-NEXT: acc.detach accPtr(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<host>]) {dataClause = #acc<data_clause acc_attach>, name = "arg1"}1187 1188 // Checking the automatic-addition of parallelism clauses.1189#pragma acc parallel loop1190 for(unsigned I = 0; I < 5; ++I);1191 // CHECK-NEXT: acc.parallel combined(loop) {1192 // CHECK-NEXT: acc.loop combined(parallel) {1193 // CHECK: acc.yield1194 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc1195 // CHECK-NEXT: acc.yield1196 // CHECK-NEXT: } loc1197 1198#pragma acc kernels loop1199 for(unsigned I = 0; I < 5; ++I);1200 // CHECK-NEXT: acc.kernels combined(loop) {1201 // CHECK-NEXT: acc.loop combined(kernels) {1202 // CHECK: acc.yield1203 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc1204 // CHECK-NEXT: acc.terminator1205 // CHECK-NEXT: } loc1206 1207#pragma acc serial loop1208 for(unsigned I = 0; I < 5; ++I);1209 // CHECK-NEXT: acc.serial combined(loop) {1210 // CHECK-NEXT: acc.loop combined(serial) {1211 // CHECK: acc.yield1212 // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc1213 // CHECK-NEXT: acc.yield1214 // CHECK-NEXT: } loc1215 1216#pragma acc serial loop worker1217 for(unsigned I = 0; I < 5; ++I);1218 // CHECK-NEXT: acc.serial combined(loop) {1219 // CHECK-NEXT: acc.loop combined(serial) worker {1220 // CHECK: acc.yield1221 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc1222 // CHECK-NEXT: acc.yield1223 // CHECK-NEXT: } loc1224 1225#pragma acc serial loop vector1226 for(unsigned I = 0; I < 5; ++I);1227 // CHECK-NEXT: acc.serial combined(loop) {1228 // CHECK-NEXT: acc.loop combined(serial) vector {1229 // CHECK: acc.yield1230 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc1231 // CHECK-NEXT: acc.yield1232 // CHECK-NEXT: } loc1233 1234#pragma acc serial loop gang1235 for(unsigned I = 0; I < 5; ++I);1236 // CHECK-NEXT: acc.serial combined(loop) {1237 // CHECK-NEXT: acc.loop combined(serial) gang {1238 // CHECK: acc.yield1239 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc1240 // CHECK-NEXT: acc.yield1241 // CHECK-NEXT: } loc1242}1243