274 lines · c
1// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s2 3void acc_data(int cond) {4 // CHECK: cir.func{{.*}} @acc_data(%[[ARG:.*]]: !s32i{{.*}}) {5 // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]6 7 int *ptr;8 // CHECK-NEXT: %[[PTR:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["ptr"]9 // CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>10 11#pragma acc data default(none)12 {13 int i = 0;14 ++i;15 }16 // CHECK-NEXT: acc.data {17 // CHECK-NEXT: cir.alloca18 // CHECK-NEXT: cir.const19 // CHECK-NEXT: cir.store20 // CHECK-NEXT: cir.load21 // CHECK-NEXT: cir.unary22 // CHECK-NEXT: cir.store23 // CHECK-NEXT: acc.terminator24 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}25 26#pragma acc data default(present)27 {28 int i = 0;29 ++i;30 }31 // CHECK-NEXT: acc.data {32 // CHECK-NEXT: cir.alloca33 // CHECK-NEXT: cir.const34 // CHECK-NEXT: cir.store35 // CHECK-NEXT: cir.load36 // CHECK-NEXT: cir.unary37 // CHECK-NEXT: cir.store38 // CHECK-NEXT: acc.terminator39 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue present>}40 41#pragma acc data default(none) async42 {}43 // CHECK-NEXT: acc.data async {44 // CHECK-NEXT: acc.terminator45 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}46 47#pragma acc data default(none) async(cond)48 {}49 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i50 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si3251 // CHECK-NEXT: acc.data async(%[[CONV_CAST]] : si32) {52 // CHECK-NEXT: acc.terminator53 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}54 55#pragma acc data default(none) async device_type(nvidia, radeon) async56 {}57 // CHECK-NEXT: acc.data async([#acc.device_type<none>, #acc.device_type<nvidia>, #acc.device_type<radeon>]) {58 // CHECK-NEXT: acc.terminator59 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}60 61#pragma acc data default(none) async(3) device_type(nvidia, radeon) async(cond)62 {}63 // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i64 // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si3265 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i66 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si3267 // CHECK-NEXT: acc.data async(%[[THREE_CAST]] : si32, %[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {68 // CHECK-NEXT: acc.terminator69 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}70 71#pragma acc data default(none) async device_type(nvidia, radeon) async(cond)72 {}73 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i74 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si3275 // CHECK-NEXT: acc.data async([#acc.device_type<none>], %[[CONV_CAST]] : si32 [#acc.device_type<nvidia>], %[[CONV_CAST]] : si32 [#acc.device_type<radeon>]) {76 // CHECK-NEXT: acc.terminator77 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}78 79#pragma acc data default(none) async(3) device_type(nvidia, radeon) async80 {}81 // CHECK-NEXT: %[[THREE_LITERAL:.*]] = cir.const #cir.int<3> : !s32i82 // CHECK-NEXT: %[[THREE_CAST:.*]] = builtin.unrealized_conversion_cast %[[THREE_LITERAL]] : !s32i to si3283 // CHECK-NEXT: acc.data async([#acc.device_type<nvidia>, #acc.device_type<radeon>], %[[THREE_CAST]] : si32) {84 // CHECK-NEXT: acc.terminator85 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}86 87#pragma acc data default(none) if(cond)88 {}89 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i90 // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast int_to_bool %[[COND_LOAD]] : !s32i -> !cir.bool91 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i192 // CHECK-NEXT: acc.data if(%[[CONV_CAST]]) {93 // CHECK-NEXT: acc.terminator94 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}95 96#pragma acc data default(none) if(1)97 {}98 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i99 // CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast int_to_bool %[[ONE_LITERAL]] : !s32i -> !cir.bool100 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1101 // CHECK-NEXT: acc.data if(%[[CONV_CAST]]) {102 // CHECK-NEXT: acc.terminator103 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}104 105#pragma acc data default(none) if(cond == 1)106 {}107 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i108 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i109 // CHECK-NEXT: %[[EQ_RES:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool110 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES]] : !cir.bool to i1111 // CHECK-NEXT: acc.data if(%[[CONV_CAST]]) {112 // CHECK-NEXT: acc.terminator113 // CHECK-NEXT: } attributes {defaultAttr = #acc<defaultvalue none>}114 115#pragma acc data default(none) wait116 {}117 // CHECK-NEXT: acc.data wait {118 // CHECK-NEXT: acc.terminator119 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}120 121#pragma acc data default(none) wait device_type(nvidia) wait122 {}123 // CHECK-NEXT: acc.data wait([#acc.device_type<none>, #acc.device_type<nvidia>]) {124 // CHECK-NEXT: acc.terminator125 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}126 127#pragma acc data default(none) wait(1) device_type(nvidia) wait128 {}129 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i130 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32131 // CHECK-NEXT: acc.data wait([#acc.device_type<nvidia>], {%[[ONE_CAST]] : si32}) {132 // CHECK-NEXT: acc.terminator133 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}134 135#pragma acc data default(none) wait device_type(nvidia) wait(1)136 {}137 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i138 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32139 // CHECK-NEXT: acc.data wait([#acc.device_type<none>], {%[[ONE_CAST]] : si32} [#acc.device_type<nvidia>]) {140 // CHECK-NEXT: acc.terminator141 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}142 143#pragma acc data default(none) wait(1) device_type(nvidia) wait(1)144 {}145 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i146 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32147 // CHECK-NEXT: %[[ONE_LITERAL2:.*]] = cir.const #cir.int<1> : !s32i148 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL2]] : !s32i to si32149 // CHECK-NEXT: acc.data wait({%[[ONE_CAST]] : si32}, {%[[ONE_CAST2]] : si32} [#acc.device_type<nvidia>]) {150 // CHECK-NEXT: acc.terminator151 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}152 153#pragma acc data default(none) wait(devnum: cond : 1)154 {}155 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i156 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32157 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i158 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32159 // CHECK-NEXT: acc.data wait({devnum: %[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32}) {160 // CHECK-NEXT: acc.terminator161 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}162 163#pragma acc data default(none) wait(devnum: cond : 1) device_type(nvidia) wait(devnum: cond : 1)164 {}165 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i166 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32167 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i168 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32169 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i170 // CHECK-NEXT: %[[CONV_CAST2:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32171 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i172 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32173 // CHECK-NEXT: acc.data wait({devnum: %[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32}, {devnum: %[[CONV_CAST2]] : si32, %[[ONE_CAST2]] : si32} [#acc.device_type<nvidia>]) {174 // CHECK-NEXT: acc.terminator175 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}176 177#pragma acc data default(none) wait(devnum: cond : 1, 2)178 {}179 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i180 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32181 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i182 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32183 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i184 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32185 // CHECK-NEXT: acc.data wait({devnum: %[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32, %[[TWO_CAST]] : si32}) {186 // CHECK-NEXT: acc.terminator187 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}188 189#pragma acc data default(none) wait(devnum: cond : 1, 2) device_type(nvidia, radeon) wait(devnum: cond : 1, 2)190 {}191 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i192 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32193 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i194 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32195 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i196 // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32197 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i198 // CHECK-NEXT: %[[CONV_CAST2:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32199 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i200 // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32201 // CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i202 // CHECK-NEXT: %[[TWO_CAST2:.*]] = builtin.unrealized_conversion_cast %[[TWO_LITERAL]] : !s32i to si32203 // CHECK-NEXT: acc.data 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>]) {204 // CHECK-NEXT: acc.terminator205 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}206 207#pragma acc data default(none) wait(cond, 1)208 {}209 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i210 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32211 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i212 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32213 // CHECK-NEXT: acc.data wait({%[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32}) {214 // CHECK-NEXT: acc.terminator215 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}216 217#pragma acc data default(none) wait(queues: cond, 1) device_type(radeon)218 {}219 // CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i220 // CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32221 // CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i222 // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32223 // CHECK-NEXT: acc.data wait({%[[CONV_CAST]] : si32, %[[ONE_CAST]] : si32}) {224 // CHECK-NEXT: acc.terminator225 // CHECK-NEXT: attributes {defaultAttr = #acc<defaultvalue none>}226 227#pragma acc data deviceptr(ptr)228 {}229 // CHECK-NEXT: %[[DEV_PTR:.*]] = acc.deviceptr varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"}230 // CHECK-NEXT: acc.data dataOperands(%[[DEV_PTR]] : !cir.ptr<!cir.ptr<!s32i>>) {231 // CHECK-NEXT: acc.terminator232 // CHECK-NEXT: } loc233#pragma acc data deviceptr(ptr) device_type(radeon) async234 {}235 // CHECK-NEXT: %[[DEV_PTR:.*]] = acc.deviceptr varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"}236 // CHECK-NEXT: acc.data async([#acc.device_type<radeon>]) dataOperands(%[[DEV_PTR]] : !cir.ptr<!cir.ptr<!s32i>>) {237 // CHECK-NEXT: acc.terminator238 // CHECK-NEXT: } loc239 240#pragma acc data present(cond)241 {}242 // CHECK-NEXT: %[[PRESENT:.*]] = acc.present varPtr(%[[COND]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "cond"}243 // CHECK-NEXT: acc.data dataOperands(%[[PRESENT]] : !cir.ptr<!s32i>) {244 // CHECK-NEXT: acc.terminator245 // CHECK-NEXT: } loc246 // CHECK-NEXT: acc.delete accPtr(%[[PRESENT]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_present>, name = "cond"}247 248#pragma acc data present(cond) device_type(radeon) async249 {}250 // CHECK-NEXT: %[[PRESENT:.*]] = acc.present varPtr(%[[COND]] : !cir.ptr<!s32i>) async([#acc.device_type<radeon>]) -> !cir.ptr<!s32i> {name = "cond"}251 // CHECK-NEXT: acc.data async([#acc.device_type<radeon>]) dataOperands(%[[PRESENT]] : !cir.ptr<!s32i>) {252 // CHECK-NEXT: acc.terminator253 // CHECK-NEXT: } loc254 // CHECK-NEXT: acc.delete accPtr(%[[PRESENT]] : !cir.ptr<!s32i>) async([#acc.device_type<radeon>]) {dataClause = #acc<data_clause acc_present>, name = "cond"}255 256#pragma acc data attach(ptr)257 {}258 // CHECK-NEXT: %[[ATTACH:.*]] = acc.attach varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"}259 // CHECK-NEXT: acc.data dataOperands(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) {260 // CHECK-NEXT: acc.terminator261 // CHECK-NEXT: } loc262 // CHECK-NEXT: acc.detach accPtr(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_attach>, name = "ptr"}263 264#pragma acc data attach(ptr) device_type(radeon) async265 {}266 // CHECK-NEXT: %[[ATTACH:.*]] = acc.attach varPtr(%[[PTR]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptr"}267 // CHECK-NEXT: acc.data async([#acc.device_type<radeon>]) dataOperands(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) {268 // CHECK-NEXT: acc.terminator269 // CHECK-NEXT: } loc270 // CHECK-NEXT: acc.detach accPtr(%[[ATTACH]] : !cir.ptr<!cir.ptr<!s32i>>) async([#acc.device_type<radeon>]) {dataClause = #acc<data_clause acc_attach>, name = "ptr"}271 272 // CHECK-NEXT: cir.return273}274