Skip to content

Commit c05da6e

Browse files
committed
[OpenACC][CIR] Implement lowering for 'set' clauses
The 'set' clauses are default_async, device_num, and if. The latter two are implemented identically to other constructs by that name. default_async works exactly like device_num, (and others) that take an int-expression.
1 parent f52b01b commit c05da6e

File tree

2 files changed

+56
-4
lines changed

2 files changed

+56
-4
lines changed

clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp

Lines changed: 24 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -170,6 +170,7 @@ class OpenACCClauseCIREmitter final
170170
break;
171171
}
172172
} else {
173+
// Combined Constructs left.
173174
return clauseNotImplemented(clause);
174175
}
175176
}
@@ -208,6 +209,7 @@ class OpenACCClauseCIREmitter final
208209
// they just modify the other clauses IR. So setting of `lastDeviceType`
209210
// (done above) is all we need.
210211
} else {
212+
// update, data, loop, routine, combined remain.
211213
return clauseNotImplemented(clause);
212214
}
213215
}
@@ -221,6 +223,7 @@ class OpenACCClauseCIREmitter final
221223
} else if constexpr (isOneOfTypes<OpTy, SerialOp>) {
222224
llvm_unreachable("num_workers not valid on serial");
223225
} else {
226+
// Combined Remain.
224227
return clauseNotImplemented(clause);
225228
}
226229
}
@@ -234,6 +237,7 @@ class OpenACCClauseCIREmitter final
234237
} else if constexpr (isOneOfTypes<OpTy, SerialOp>) {
235238
llvm_unreachable("vector_length not valid on serial");
236239
} else {
240+
// Combined remain.
237241
return clauseNotImplemented(clause);
238242
}
239243
}
@@ -250,6 +254,7 @@ class OpenACCClauseCIREmitter final
250254
createIntExpr(clause.getIntExpr()), &range));
251255
}
252256
} else {
257+
// Data, enter data, exit data, update, wait, combined remain.
253258
return clauseNotImplemented(clause);
254259
}
255260
}
@@ -266,19 +271,21 @@ class OpenACCClauseCIREmitter final
266271
llvm_unreachable("var-list version of self shouldn't get here");
267272
}
268273
} else {
274+
// update and combined remain.
269275
return clauseNotImplemented(clause);
270276
}
271277
}
272278

273279
void VisitIfClause(const OpenACCIfClause &clause) {
274280
if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, InitOp,
275-
ShutdownOp>) {
281+
ShutdownOp, SetOp>) {
276282
operation.getIfCondMutable().append(
277283
createCondition(clause.getConditionExpr()));
278284
} else {
279285
// 'if' applies to most of the constructs, but hold off on lowering them
280286
// until we can write tests/know what we're doing with codegen to make
281287
// sure we get it right.
288+
// Enter data, exit data, host_data, update, wait, combined remain.
282289
return clauseNotImplemented(clause);
283290
}
284291
}
@@ -287,8 +294,23 @@ class OpenACCClauseCIREmitter final
287294
if constexpr (isOneOfTypes<OpTy, InitOp, ShutdownOp>) {
288295
operation.getDeviceNumOperandMutable().append(
289296
createIntExpr(clause.getIntExpr()));
297+
} else if constexpr (isOneOfTypes<OpTy, SetOp>) {
298+
// This is only a separate case because the getter name is different in
299+
// 'set' for some reason.
300+
operation.getDeviceNumMutable().append(
301+
createIntExpr(clause.getIntExpr()));
290302
} else {
291-
return clauseNotImplemented(clause);
303+
llvm_unreachable(
304+
"init, shutdown, set, are only valid device_num constructs");
305+
}
306+
}
307+
308+
void VisitDefaultAsyncClause(const OpenACCDefaultAsyncClause &clause) {
309+
if constexpr (isOneOfTypes<OpTy, SetOp>) {
310+
operation.getDefaultAsyncMutable().append(
311+
createIntExpr(clause.getIntExpr()));
312+
} else {
313+
llvm_unreachable("set, is only valid device_num constructs");
292314
}
293315
}
294316
};

clang/test/CIR/CodeGenOpenACC/set.c

Lines changed: 32 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,9 @@
11
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
22

3-
void acc_set(void) {
4-
// CHECK: cir.func @acc_set() {
3+
void acc_set(int cond) {
4+
// CHECK: cir.func @acc_set(%[[ARG:.*]]: !s32i{{.*}}) {
5+
// CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
6+
// CHECK-NEXT: cir.store %[[ARG]], %[[COND]] : !s32i, !cir.ptr<!s32i>
57

68
#pragma acc set device_type(*)
79
// CHECK-NEXT: acc.set attributes {device_type = #acc.device_type<star>}
@@ -10,5 +12,33 @@ void acc_set(void) {
1012
#pragma acc set device_type(radeon)
1113
// CHECK-NEXT: acc.set attributes {device_type = #acc.device_type<radeon>}
1214

15+
#pragma acc set default_async(cond)
16+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
17+
// CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
18+
// CHECK-NEXT: acc.set default_async(%[[COND_CONV]] : si32)
19+
20+
#pragma acc set default_async(1)
21+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
22+
// CHECK-NEXT: %[[ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
23+
// CHECK-NEXT: acc.set default_async(%[[ONE_CONV]] : si32)
24+
25+
#pragma acc set device_num(cond) if (cond)
26+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
27+
// CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
28+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
29+
// CHECK-NEXT: %[[COND_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
30+
// CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_CAST]] : !cir.bool to i1
31+
// CHECK-NEXT: acc.set device_num(%[[COND_CONV]] : si32) if(%[[BOOL_CONV]])
32+
33+
#pragma acc set device_type(radeon) default_async(1) device_num(cond) if (cond)
34+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
35+
// CHECK-NEXT: %[[ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
36+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
37+
// CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
38+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
39+
// CHECK-NEXT: %[[COND_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
40+
// CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_CAST]] : !cir.bool to i1
41+
// CHECK-NEXT: acc.set default_async(%[[ONE_CONV]] : si32) device_num(%[[COND_CONV]] : si32) if(%[[BOOL_CONV]]) attributes {device_type = #acc.device_type<radeon>}
42+
1343
// CHECK-NEXT: cir.return
1444
}

0 commit comments

Comments
 (0)