Skip to content

Commit ab4e181

Browse files
erichkeaneerichkeane
committed
[OpenACC][CIR] Implement 'if' and 'device_num' lowering for
init/shutdown These are pretty simple, the 'if' implementation is the same as compute constructs, and the 'device_num' is identical to a bunch of others, in that it is just emitting an integral value.
1 parent 0eba8cb commit ab4e181

File tree

3 files changed

+79
-5
lines changed

3 files changed

+79
-5
lines changed

‎clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp

Lines changed: 11 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -271,7 +271,8 @@ class OpenACCClauseCIREmitter final
271271
}
272272

273273
void VisitIfClause(const OpenACCIfClause &clause) {
274-
if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
274+
if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp, InitOp,
275+
ShutdownOp>) {
275276
operation.getIfCondMutable().append(
276277
createCondition(clause.getConditionExpr()));
277278
} else {
@@ -281,6 +282,15 @@ class OpenACCClauseCIREmitter final
281282
return clauseNotImplemented(clause);
282283
}
283284
}
285+
286+
void VisitDeviceNumClause(const OpenACCDeviceNumClause &clause) {
287+
if constexpr (isOneOfTypes<OpTy, InitOp, ShutdownOp>) {
288+
operation.getDeviceNumOperandMutable().append(
289+
createIntExpr(clause.getIntExpr()));
290+
} else {
291+
return clauseNotImplemented(clause);
292+
}
293+
}
284294
};
285295

286296
template <typename OpTy>

‎clang/test/CIR/CodeGenOpenACC/init.c

Lines changed: 34 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_init(void) {
4-
// CHECK: cir.func @acc_init() {
3+
void acc_init(int cond) {
4+
// CHECK: cir.func @acc_init(%[[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
#pragma acc init
68
// CHECK-NEXT: acc.init loc(#{{[a-zA-Z0-9]+}}){{$}}
79

@@ -17,4 +19,34 @@ void acc_init(void) {
1719
// CHECK-NEXT: acc.init attributes {device_types = [#acc.device_type<host>, #acc.device_type<multicore>]}
1820
#pragma acc init device_type(HoSt) device_type(MuLtIcORe)
1921
// CHECK-NEXT: acc.init attributes {device_types = [#acc.device_type<host>, #acc.device_type<multicore>]}
22+
23+
#pragma acc init if(cond)
24+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
25+
// CHECK-NEXT: %[[COND_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
26+
// CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_CAST]] : !cir.bool to i1
27+
// CHECK-NEXT: acc.init if(%[[BOOL_CONV]])
28+
29+
#pragma acc init if(1)
30+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
31+
// CHECK-NEXT: %[[ONE_TO_BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ONE_LITERAL]] : !s32i), !cir.bool
32+
// CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[ONE_TO_BOOL_CAST]] : !cir.bool to i1
33+
// CHECK-NEXT: acc.init if(%[[BOOL_CONV]])
34+
35+
#pragma acc init device_num(cond)
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: acc.init device_num(%[[COND_CONV]] : si32)
39+
40+
#pragma acc init device_num(1)
41+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
42+
// CHECK-NEXT: %[[ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
43+
// CHECK-NEXT: acc.init device_num(%[[ONE_CONV]] : si32)
44+
45+
#pragma acc init if(cond) device_num(cond) device_type(*)
46+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
47+
// CHECK-NEXT: %[[COND_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
48+
// CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_CAST]] : !cir.bool to i1
49+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
50+
// CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
51+
// CHECK-NEXT: acc.init device_num(%[[COND_CONV]] : si32) if(%[[BOOL_CONV]]) attributes {device_types = [#acc.device_type<star>]}
2052
}

‎clang/test/CIR/CodeGenOpenACC/shutdown.c

Lines changed: 34 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_shutdown(void) {
4-
// CHECK: cir.func @acc_shutdown() {
3+
void acc_shutdown(int cond) {
4+
// CHECK: cir.func @acc_shutdown(%[[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
#pragma acc shutdown
68
// CHECK-NEXT: acc.shutdown loc(#{{[a-zA-Z0-9]+}}){{$}}
79

@@ -17,4 +19,34 @@ void acc_shutdown(void) {
1719
// CHECK-NEXT: acc.shutdown attributes {device_types = [#acc.device_type<host>, #acc.device_type<multicore>]}
1820
#pragma acc shutdown device_type(HoSt) device_type(MuLtIcORe)
1921
// CHECK-NEXT: acc.shutdown attributes {device_types = [#acc.device_type<host>, #acc.device_type<multicore>]}
22+
23+
#pragma acc shutdown if(cond)
24+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
25+
// CHECK-NEXT: %[[COND_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
26+
// CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_CAST]] : !cir.bool to i1
27+
// CHECK-NEXT: acc.shutdown if(%[[BOOL_CONV]])
28+
29+
#pragma acc shutdown if(1)
30+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
31+
// CHECK-NEXT: %[[ONE_TO_BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ONE_LITERAL]] : !s32i), !cir.bool
32+
// CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[ONE_TO_BOOL_CAST]] : !cir.bool to i1
33+
// CHECK-NEXT: acc.shutdown if(%[[BOOL_CONV]])
34+
35+
#pragma acc shutdown device_num(cond)
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: acc.shutdown device_num(%[[COND_CONV]] : si32)
39+
40+
#pragma acc shutdown device_num(1)
41+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
42+
// CHECK-NEXT: %[[ONE_CONV:.*]] = builtin.unrealized_conversion_cast %[[ONE_LITERAL]] : !s32i to si32
43+
// CHECK-NEXT: acc.shutdown device_num(%[[ONE_CONV]] : si32)
44+
45+
#pragma acc shutdown if(cond) device_num(cond) device_type(*)
46+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
47+
// CHECK-NEXT: %[[COND_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
48+
// CHECK-NEXT: %[[BOOL_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_CAST]] : !cir.bool to i1
49+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
50+
// CHECK-NEXT: %[[COND_CONV:.*]] = builtin.unrealized_conversion_cast %[[COND_LOAD]] : !s32i to si32
51+
// CHECK-NEXT: acc.shutdown device_num(%[[COND_CONV]] : si32) if(%[[BOOL_CONV]]) attributes {device_types = [#acc.device_type<star>]}
2052
}

0 commit comments

Comments
 (0)