477 lines · cpp
1// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s2 3extern "C" void acc_loop(int *A, int *B, int *C, int N) {4 // CHECK: cir.func{{.*}} @acc_loop(%[[ARG_A:.*]]: !cir.ptr<!s32i> loc{{.*}}, %[[ARG_B:.*]]: !cir.ptr<!s32i> loc{{.*}}, %[[ARG_C:.*]]: !cir.ptr<!s32i> loc{{.*}}, %[[ARG_N:.*]]: !s32i loc{{.*}}) {5 // CHECK-NEXT: %[[ALLOCA_A:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["A", init]6 // CHECK-NEXT: %[[ALLOCA_B:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["B", init]7 // CHECK-NEXT: %[[ALLOCA_C:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["C", init]8 // CHECK-NEXT: %[[ALLOCA_N:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["N", init]9 // CHECK-NEXT: cir.store %[[ARG_A]], %[[ALLOCA_A]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>10 // CHECK-NEXT: cir.store %[[ARG_B]], %[[ALLOCA_B]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>11 // CHECK-NEXT: cir.store %[[ARG_C]], %[[ALLOCA_C]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>12 // CHECK-NEXT: cir.store %[[ARG_N]], %[[ALLOCA_N]] : !s32i, !cir.ptr<!s32i>13 14 15#pragma acc loop16 for (unsigned I = 0u; I < N; ++I) {17 A[I] = B[I] + C[I];18 }19 // CHECK-NEXT: acc.loop {20 // CHECK-NEXT: cir.scope {21 // CHECK: cir.for : cond {22 // CHECK: cir.condition23 // CHECK-NEXT: } body {24 // CHECK-NEXT: cir.scope {25 // CHECK: }26 // CHECK-NEXT: cir.yield27 // CHECK-NEXT: } step {28 // CHECK: cir.yield29 // CHECK-NEXT: } loc30 // CHECK-NEXT: } loc31 // CHECK-NEXT: acc.yield32 // CHECK-NEXT: } loc33 34 35#pragma acc loop seq36 for(unsigned I = 0; I < N; ++I);37 // CHECK: acc.loop {38 // CHECK: acc.yield39 // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc40#pragma acc loop device_type(nvidia, radeon) seq41 for(unsigned I = 0; I < N; ++I);42 // CHECK: acc.loop {43 // CHECK: acc.yield44 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>], seq = [#acc.device_type<nvidia>, #acc.device_type<radeon>]} loc45#pragma acc loop device_type(radeon) seq46 for(unsigned I = 0; I < N; ++I);47 // CHECK: acc.loop {48 // CHECK: acc.yield49 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>], seq = [#acc.device_type<radeon>]} loc50#pragma acc loop seq device_type(nvidia, radeon)51 for(unsigned I = 0; I < N; ++I);52 // CHECK: acc.loop {53 // CHECK: acc.yield54 // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc55#pragma acc loop seq device_type(radeon)56 for(unsigned I = 0; I < N; ++I);57 // CHECK: acc.loop {58 // CHECK: acc.yield59 // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc60 61#pragma acc loop independent62 for(unsigned I = 0; I < N; ++I);63 // CHECK: acc.loop {64 // CHECK: acc.yield65 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc66#pragma acc loop device_type(nvidia, radeon) independent67 for(unsigned I = 0; I < N; ++I);68 // CHECK: acc.loop {69 // CHECK: acc.yield70 // CHECK-NEXT: } attributes {independent = [#acc.device_type<nvidia>, #acc.device_type<radeon>, #acc.device_type<none>]} loc71#pragma acc loop device_type(radeon) independent72 for(unsigned I = 0; I < N; ++I);73 // CHECK: acc.loop {74 // CHECK: acc.yield75 // CHECK-NEXT: } attributes {independent = [#acc.device_type<radeon>, #acc.device_type<none>]} loc76#pragma acc loop independent device_type(nvidia, radeon)77 for(unsigned I = 0; I < N; ++I);78 // CHECK: acc.loop {79 // CHECK: acc.yield80 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc81#pragma acc loop independent device_type(radeon)82 for(unsigned I = 0; I < N; ++I);83 // CHECK: acc.loop {84 // CHECK: acc.yield85 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc86 87#pragma acc loop auto88 for(unsigned I = 0; I < N; ++I);89 // CHECK: acc.loop {90 // CHECK: acc.yield91 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc92#pragma acc loop device_type(nvidia, radeon) auto93 for(unsigned I = 0; I < N; ++I);94 // CHECK: acc.loop {95 // CHECK: acc.yield96 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<nvidia>, #acc.device_type<radeon>], independent = [#acc.device_type<none>]} loc97#pragma acc loop device_type(radeon) auto98 for(unsigned I = 0; I < N; ++I);99 // CHECK: acc.loop {100 // CHECK: acc.yield101 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<radeon>], independent = [#acc.device_type<none>]} loc102#pragma acc loop auto device_type(nvidia, radeon)103 for(unsigned I = 0; I < N; ++I);104 // CHECK: acc.loop {105 // CHECK: acc.yield106 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc107#pragma acc loop auto device_type(radeon)108 for(unsigned I = 0; I < N; ++I);109 // CHECK: acc.loop {110 // CHECK: acc.yield111 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc112 113 #pragma acc loop collapse(1) device_type(radeon)114 for(unsigned I = 0; I < N; ++I)115 for(unsigned J = 0; J < N; ++J)116 for(unsigned K = 0; K < N; ++K);117 // CHECK: acc.loop {118 // CHECK: acc.yield119 // CHECK-NEXT: } attributes {collapse = [1], collapseDeviceType = [#acc.device_type<none>], independent = [#acc.device_type<none>]}120 121 #pragma acc loop collapse(1) device_type(radeon) collapse (2)122 for(unsigned I = 0; I < N; ++I)123 for(unsigned J = 0; J < N; ++J)124 for(unsigned K = 0; K < N; ++K);125 // CHECK: acc.loop {126 // CHECK: acc.yield127 // CHECK-NEXT: } attributes {collapse = [1, 2], collapseDeviceType = [#acc.device_type<none>, #acc.device_type<radeon>], independent = [#acc.device_type<none>]}128 129 #pragma acc loop collapse(1) device_type(radeon, nvidia) collapse (2)130 for(unsigned I = 0; I < N; ++I)131 for(unsigned J = 0; J < N; ++J)132 for(unsigned K = 0; K < N; ++K);133 // CHECK: acc.loop {134 // CHECK: acc.yield135 // CHECK-NEXT: } attributes {collapse = [1, 2, 2], collapseDeviceType = [#acc.device_type<none>, #acc.device_type<radeon>, #acc.device_type<nvidia>], independent = [#acc.device_type<none>]}136 #pragma acc loop collapse(1) device_type(radeon, nvidia) collapse(2) device_type(host) collapse(3)137 for(unsigned I = 0; I < N; ++I)138 for(unsigned J = 0; J < N; ++J)139 for(unsigned K = 0; K < N; ++K);140 // CHECK: acc.loop {141 // CHECK: acc.yield142 // 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>]}143 144 #pragma acc loop tile(1, 2, 3)145 for(unsigned I = 0; I < N; ++I)146 for(unsigned J = 0; J < N; ++J)147 for(unsigned K = 0; K < N; ++K);148 // CHECK: %[[ONE_CONST:.*]] = arith.constant 1 : i64149 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64150 // CHECK-NEXT: %[[THREE_CONST:.*]] = arith.constant 3 : i64151 // CHECK-NEXT: acc.loop tile({%[[ONE_CONST]] : i64, %[[TWO_CONST]] : i64, %[[THREE_CONST]] : i64}) {152 // CHECK: acc.yield153 // CHECK-NEXT: } loc154 #pragma acc loop tile(2) device_type(radeon)155 for(unsigned I = 0; I < N; ++I)156 for(unsigned J = 0; J < N; ++J)157 for(unsigned K = 0; K < N; ++K);158 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64159 // CHECK-NEXT: acc.loop tile({%[[TWO_CONST]] : i64}) {160 // CHECK: acc.yield161 // CHECK-NEXT: } loc162 #pragma acc loop tile(2) device_type(radeon) tile (1, *)163 for(unsigned I = 0; I < N; ++I)164 for(unsigned J = 0; J < N; ++J)165 for(unsigned K = 0; K < N; ++K);166 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64167 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64168 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64169 // CHECK-NEXT: acc.loop tile({%[[TWO_CONST]] : i64}, {%[[ONE_CONST]] : i64, %[[STAR_CONST]] : i64} [#acc.device_type<radeon>]) {170 // CHECK: acc.yield171 // CHECK-NEXT: } loc172 #pragma acc loop tile(*) device_type(radeon, nvidia) tile (1, 2)173 for(unsigned I = 0; I < N; ++I)174 for(unsigned J = 0; J < N; ++J)175 for(unsigned K = 0; K < N; ++K);176 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64177 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64178 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64179 // CHECK-NEXT: acc.loop tile({%[[STAR_CONST]] : i64}, {%[[ONE_CONST]] : i64, %[[TWO_CONST]] : i64} [#acc.device_type<radeon>], {%[[ONE_CONST]] : i64, %[[TWO_CONST]] : i64} [#acc.device_type<nvidia>]) {180 // CHECK: acc.yield181 // CHECK-NEXT: } loc182 #pragma acc loop tile(1) device_type(radeon, nvidia) tile(2, 3) device_type(host) tile(*, *, *)183 for(unsigned I = 0; I < N; ++I)184 for(unsigned J = 0; J < N; ++J)185 for(unsigned K = 0; K < N; ++K);186 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64187 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64188 // CHECK-NEXT: %[[THREE_CONST:.*]] = arith.constant 3 : i64189 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64190 // CHECK-NEXT: %[[STAR2_CONST:.*]] = arith.constant -1 : i64191 // CHECK-NEXT: %[[STAR3_CONST:.*]] = arith.constant -1 : i64192 // CHECK-NEXT: acc.loop 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>]) {193 // CHECK: acc.yield194 // CHECK-NEXT: } loc195 196 197#pragma acc kernels198 {199 200#pragma acc loop worker201 for(unsigned I = 0; I < N; ++I);202 // CHECK: acc.loop worker {203 // CHECK: acc.yield204 // CHECK-NEXT: } loc205 206#pragma acc loop worker(N)207 for(unsigned I = 0; I < N; ++I);208 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i209 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32210 // CHECK-NEXT: acc.loop worker(%[[N_CONV]] : si32) {211 // CHECK: acc.yield212 // CHECK-NEXT: } loc213 214#pragma acc loop worker device_type(nvidia, radeon) worker215 for(unsigned I = 0; I < N; ++I);216 // CHECK-NEXT: acc.loop worker([#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>]) {217 // CHECK: acc.yield218 // CHECK-NEXT: } loc219 220#pragma acc loop worker(N) device_type(nvidia, radeon) worker221 for(unsigned I = 0; I < N; ++I);222 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i223 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32224 // CHECK-NEXT: acc.loop worker([#acc.device_type<nvidia>, #acc.device_type<radeon>], %[[N_CONV]] : si32) {225 // CHECK: acc.yield226 // CHECK-NEXT: } loc227 228#pragma acc loop worker device_type(nvidia, radeon) worker(N)229 for(unsigned I = 0; I < N; ++I);230 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i231 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32232 // CHECK-NEXT: acc.loop worker([#acc.device_type<none>], %[[N_CONV]] : si32 [#acc.device_type<nvidia>], %[[N_CONV]] : si32 [#acc.device_type<radeon>]) {233 // CHECK: acc.yield234 // CHECK-NEXT: } loc235 236#pragma acc loop worker(N) device_type(nvidia, radeon) worker(N + 1)237 for(unsigned I = 0; I < N; ++I);238 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i239 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32240 // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i241 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i242 // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD2]], %[[ONE_CONST]]) nsw : !s32i243 // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32244 // CHECK-NEXT: acc.loop worker(%[[N_CONV]] : si32, %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<nvidia>], %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<radeon>]) {245 // CHECK: acc.yield246 // CHECK-NEXT: } loc247 248#pragma acc loop device_type(nvidia, radeon) worker(num:N + 1)249 for(unsigned I = 0; I < N; ++I);250 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i251 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i252 // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD]], %[[ONE_CONST]]) nsw : !s32i253 // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32254 // CHECK-NEXT: acc.loop worker(%[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<nvidia>], %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<radeon>]) {255 256#pragma acc loop vector257 for(unsigned I = 0; I < N; ++I);258 // CHECK: acc.loop vector {259 // CHECK: acc.yield260 // CHECK-NEXT: } loc261 262#pragma acc loop vector(N)263 for(unsigned I = 0; I < N; ++I);264 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i265 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32266 // CHECK-NEXT: acc.loop vector(%[[N_CONV]] : si32) {267 // CHECK: acc.yield268 // CHECK-NEXT: } loc269 270#pragma acc loop vector device_type(nvidia, radeon) vector271 for(unsigned I = 0; I < N; ++I);272 // CHECK-NEXT: acc.loop vector([#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>]) {273 // CHECK: acc.yield274 // CHECK-NEXT: } loc275 276#pragma acc loop vector(N) device_type(nvidia, radeon) vector277 for(unsigned I = 0; I < N; ++I);278 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i279 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32280 // CHECK-NEXT: acc.loop vector([#acc.device_type<nvidia>, #acc.device_type<radeon>], %[[N_CONV]] : si32) {281 // CHECK: acc.yield282 // CHECK-NEXT: } loc283 284#pragma acc loop vector(N) device_type(nvidia, radeon) vector(N + 1)285 for(unsigned I = 0; I < N; ++I);286 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i287 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32288 // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i289 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i290 // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD2]], %[[ONE_CONST]]) nsw : !s32i291 // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32292 // CHECK-NEXT: acc.loop vector(%[[N_CONV]] : si32, %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<nvidia>], %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<radeon>]) {293 // CHECK: acc.yield294 // CHECK-NEXT: } loc295 296#pragma acc loop device_type(nvidia, radeon) vector(length:N + 1)297 for(unsigned I = 0; I < N; ++I);298 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i299 // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i300 // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD]], %[[ONE_CONST]]) nsw : !s32i301 // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32302 // CHECK-NEXT: acc.loop vector(%[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<nvidia>], %[[N_PLUS_ONE_CONV]] : si32 [#acc.device_type<radeon>]) {303 // CHECK: acc.yield304 // CHECK-NEXT: } loc305 306#pragma acc loop worker vector device_type(nvidia) worker vector307 for(unsigned I = 0; I < N; ++I);308 // CHECK-NEXT: acc.loop worker([#acc.device_type<none>, #acc.device_type<nvidia>]) vector([#acc.device_type<none>, #acc.device_type<nvidia>])309 // CHECK: acc.yield310 // CHECK-NEXT: } loc311 312#pragma acc loop worker(N) vector(N) device_type(nvidia) worker(N) vector(N)313 for(unsigned I = 0; I < N; ++I);314 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i315 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32316 // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i317 // CHECK-NEXT: %[[N_CONV2:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD2]] : !s32i to si32318 // CHECK-NEXT: %[[N_LOAD3:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i319 // CHECK-NEXT: %[[N_CONV3:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD3]] : !s32i to si32320 // CHECK-NEXT: %[[N_LOAD4:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i321 // CHECK-NEXT: %[[N_CONV4:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD4]] : !s32i to si32322 // CHECK-NEXT: acc.loop worker(%[[N_CONV]] : si32, %[[N_CONV3]] : si32 [#acc.device_type<nvidia>]) vector(%[[N_CONV2]] : si32, %[[N_CONV4]] : si32 [#acc.device_type<nvidia>]) {323 // CHECK: acc.yield324 // CHECK-NEXT: } loc325 }326 327#pragma acc parallel328 // CHECK: acc.parallel {329 {330#pragma acc loop gang331 for(unsigned I = 0; I < N; ++I);332 // CHECK-NEXT: acc.loop gang {333 // CHECK: acc.yield334 // CHECK-NEXT: } loc335#pragma acc loop gang device_type(nvidia) gang336 for(unsigned I = 0; I < N; ++I);337 // CHECK-NEXT: acc.loop gang([#acc.device_type<none>, #acc.device_type<nvidia>]) {338 // CHECK: acc.yield339 // CHECK-NEXT: } loc340#pragma acc loop gang(dim:1) device_type(nvidia) gang(dim:2)341 for(unsigned I = 0; I < N; ++I);342 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64343 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64344 // CHECK-NEXT: acc.loop gang({dim=%[[ONE_CONST]] : i64}, {dim=%[[TWO_CONST]] : i64} [#acc.device_type<nvidia>]) {345 // CHECK: acc.yield346 // CHECK-NEXT: } loc347#pragma acc loop gang(static:N, dim: 1) device_type(nvidia, radeon) gang(static:*, dim : 2)348 for(unsigned I = 0; I < N; ++I);349 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i350 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32351 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64352 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64353 // CHECK-NEXT: %[[TWO_CONST:.*]] = arith.constant 2 : i64354 // CHECK-NEXT: acc.loop 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>]) {355 // CHECK: acc.yield356 // CHECK-NEXT: } loc357 }358#pragma acc kernels359 // CHECK: acc.kernels {360 {361#pragma acc loop gang(num:N) device_type(nvidia, radeon) gang(num:N)362 for(unsigned I = 0; I < N; ++I);363 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i364 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32365 // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i366 // CHECK-NEXT: %[[N_CONV2:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD2]] : !s32i to si32367 // CHECK-NEXT: acc.loop gang({num=%[[N_CONV]] : si32}, {num=%[[N_CONV2]] : si32} [#acc.device_type<nvidia>], {num=%[[N_CONV2]] : si32} [#acc.device_type<radeon>]) {368 // CHECK: acc.yield369 // CHECK-NEXT: } loc370#pragma acc loop gang(static:N) device_type(nvidia) gang(static:*)371 for(unsigned I = 0; I < N; ++I);372 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i373 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32374 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64375 // CHECK-NEXT: acc.loop gang({static=%[[N_CONV]] : si32}, {static=%[[STAR_CONST]] : i64} [#acc.device_type<nvidia>]) {376 // CHECK: acc.yield377 // CHECK-NEXT: } loc378#pragma acc loop gang(static:N, num: N + 1) device_type(nvidia) gang(static:*, num : N + 2)379 for(unsigned I = 0; I < N; ++I);380 // CHECK-NEXT: %[[N_LOAD:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i381 // CHECK-NEXT: %[[N_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_LOAD]] : !s32i to si32382 // CHECK-NEXT: %[[N_LOAD2:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i383 // CHECK-NEXT: %[[CIR_ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i384 // CHECK-NEXT: %[[N_PLUS_ONE:.*]] = cir.binop(add, %[[N_LOAD2]], %[[CIR_ONE_CONST]]) nsw : !s32i385 // CHECK-NEXT: %[[N_PLUS_ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_ONE]] : !s32i to si32386 // CHECK-NEXT: %[[STAR_CONST:.*]] = arith.constant -1 : i64387 // CHECK-NEXT: %[[N_LOAD3:.*]] = cir.load{{.*}} %[[ALLOCA_N]] : !cir.ptr<!s32i>, !s32i388 // CHECK-NEXT: %[[CIR_TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i389 // CHECK-NEXT: %[[N_PLUS_TWO:.*]] = cir.binop(add, %[[N_LOAD3]], %[[CIR_TWO_CONST]]) nsw : !s32i390 // CHECK-NEXT: %[[N_PLUS_TWO_CONV:.*]] = builtin.unrealized_conversion_cast %[[N_PLUS_TWO]] : !s32i to si32391 // CHECK-NEXT: acc.loop 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>]) {392 // CHECK: acc.yield393 // CHECK-NEXT: } loc394 }395 // CHECK-NEXT: acc.terminator396 // CHECK-NEXT: } loc397 398 // Checking the automatic-addition of parallelism clauses.399#pragma acc loop400 for(unsigned I = 0; I < N; ++I);401 // CHECK-NEXT: acc.loop {402 // CHECK: acc.yield403 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc404 405#pragma acc parallel406 {407 // CHECK-NEXT: acc.parallel {408#pragma acc loop409 for(unsigned I = 0; I < N; ++I);410 // CHECK-NEXT: acc.loop {411 // CHECK: acc.yield412 // CHECK-NEXT: } attributes {independent = [#acc.device_type<none>]} loc413 }414 // CHECK-NEXT: acc.yield415 // CHECK-NEXT: } loc416 417#pragma acc kernels418 {419 // CHECK-NEXT: acc.kernels {420#pragma acc loop421 for(unsigned I = 0; I < N; ++I);422 // CHECK-NEXT: acc.loop {423 // CHECK: acc.yield424 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc425 }426 // CHECK-NEXT: acc.terminator427 // CHECK-NEXT: } loc428 429#pragma acc serial430 {431 // CHECK-NEXT: acc.serial {432#pragma acc loop433 for(unsigned I = 0; I < N; ++I);434 // CHECK-NEXT: acc.loop {435 // CHECK: acc.yield436 // CHECK-NEXT: } attributes {seq = [#acc.device_type<none>]} loc437 }438 // CHECK-NEXT: acc.yield439 // CHECK-NEXT: } loc440 441#pragma acc serial442 {443 // CHECK-NEXT: acc.serial {444#pragma acc loop worker445 for(unsigned I = 0; I < N; ++I);446 // CHECK-NEXT: acc.loop worker {447 // CHECK: acc.yield448 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc449 }450 // CHECK-NEXT: acc.yield451 // CHECK-NEXT: } loc452 453#pragma acc serial454 {455 // CHECK-NEXT: acc.serial {456#pragma acc loop vector457 for(unsigned I = 0; I < N; ++I);458 // CHECK-NEXT: acc.loop vector {459 // CHECK: acc.yield460 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc461 }462 // CHECK-NEXT: acc.yield463 // CHECK-NEXT: } loc464 465#pragma acc serial466 {467 // CHECK-NEXT: acc.serial {468#pragma acc loop gang469 for(unsigned I = 0; I < N; ++I);470 // CHECK-NEXT: acc.loop gang {471 // CHECK: acc.yield472 // CHECK-NEXT: } attributes {auto_ = [#acc.device_type<none>]} loc473 }474 // CHECK-NEXT: acc.yield475 // CHECK-NEXT: } loc476}477