Skip to content

Commit 6ad922b

Browse files
erichkeaneerichkeane
committed
[OpenACC][CIR] Implement lowering for 'if' on compute constructs
This is the same for these as the 'self' was, except it doesn't support the 'empty' variant, so we have to just generate the condition. This patch does that, and extracts the 'condition' emission to a separate function since the two share it.
1 parent 860d038 commit 6ad922b

File tree

4 files changed

+159
-15
lines changed

4 files changed

+159
-15
lines changed

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

Lines changed: 29 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -50,6 +50,21 @@ class OpenACCClauseCIREmitter final
5050
cgf.cgm.errorNYI(c.getSourceRange(), "OpenACC Clause", c.getClauseKind());
5151
}
5252

53+
// 'condition' as an OpenACC grammar production is used for 'if' and (some
54+
// variants of) 'self'. It needs to be emitted as a signless-1-bit value, so
55+
// this function emits the expression, then sets the unrealized conversion
56+
// cast correctly, and returns the completed value.
57+
mlir::Value createCondition(const Expr *condExpr) {
58+
mlir::Value condition = cgf.evaluateExprAsBool(condExpr);
59+
mlir::Location exprLoc = cgf.cgm.getLoc(condExpr->getBeginLoc());
60+
mlir::IntegerType targetType = mlir::IntegerType::get(
61+
&cgf.getMLIRContext(), /*width=*/1,
62+
mlir::IntegerType::SignednessSemantics::Signless);
63+
auto conversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
64+
exprLoc, targetType, condition);
65+
return conversionOp.getResult(0);
66+
}
67+
5368
public:
5469
OpenACCClauseCIREmitter(OpTy &operation, CIRGenFunction &cgf,
5570
CIRGenBuilderTy &builder,
@@ -132,24 +147,27 @@ class OpenACCClauseCIREmitter final
132147
operation.setSelfAttr(true);
133148
} else if (clause.isConditionExprClause()) {
134149
assert(clause.hasConditionExpr());
135-
mlir::Value condition =
136-
cgf.evaluateExprAsBool(clause.getConditionExpr());
137-
138-
mlir::Location exprLoc =
139-
cgf.cgm.getLoc(clause.getConditionExpr()->getBeginLoc());
140-
mlir::IntegerType targetType = mlir::IntegerType::get(
141-
&cgf.getMLIRContext(), /*width=*/1,
142-
mlir::IntegerType::SignednessSemantics::Signless);
143-
auto conversionOp = builder.create<mlir::UnrealizedConversionCastOp>(
144-
exprLoc, targetType, condition);
145-
operation.getSelfCondMutable().append(conversionOp.getResult(0));
150+
operation.getSelfCondMutable().append(
151+
createCondition(clause.getConditionExpr()));
146152
} else {
147153
llvm_unreachable("var-list version of self shouldn't get here");
148154
}
149155
} else {
150156
return clauseNotImplemented(clause);
151157
}
152158
}
159+
160+
void VisitIfClause(const OpenACCIfClause &clause) {
161+
if constexpr (isOneOfTypes<OpTy, ParallelOp, SerialOp, KernelsOp>) {
162+
operation.getIfCondMutable().append(
163+
createCondition(clause.getConditionExpr()));
164+
} else {
165+
// 'if' applies to most of the constructs, but hold off on lowering them
166+
// until we can write tests/know what we're doing with codegen to make
167+
// sure we get it right.
168+
return clauseNotImplemented(clause);
169+
}
170+
}
153171
};
154172

155173
template <typename OpTy>

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

Lines changed: 44 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
1+
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
22

33
void acc_kernels(int cond) {
44
// CHECK: cir.func @acc_kernels(%[[ARG:.*]]: !s32i{{.*}}) {
@@ -63,6 +63,48 @@ void acc_kernels(int cond) {
6363
// CHECK-NEXT: acc.kernels self(%[[CONV_CAST]]) {
6464
// CHECK-NEXT: acc.terminator
6565
// CHECK-NEXT: } loc
66-
66+
67+
#pragma acc kernels if(cond)
68+
{}
69+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
70+
// CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
71+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
72+
// CHECK-NEXT: acc.kernels if(%[[CONV_CAST]]) {
73+
// CHECK-NEXT: acc.terminator
74+
// CHECK-NEXT: } loc
75+
76+
#pragma acc kernels if(1)
77+
{}
78+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
79+
// CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ONE_LITERAL]] : !s32i), !cir.bool
80+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
81+
// CHECK-NEXT: acc.kernels if(%[[CONV_CAST]]) {
82+
// CHECK-NEXT: acc.terminator
83+
// CHECK-NEXT: } loc
84+
85+
#pragma acc kernels if(cond == 1)
86+
{}
87+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
88+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
89+
// CHECK-NEXT: %[[EQ_RES:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool
90+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES]] : !cir.bool to i1
91+
// CHECK-NEXT: acc.kernels if(%[[CONV_CAST]]) {
92+
// CHECK-NEXT: acc.terminator
93+
// CHECK-NEXT: } loc
94+
95+
#pragma acc kernels if(cond == 1) self(cond == 2)
96+
{}
97+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
98+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
99+
// CHECK-NEXT: %[[EQ_RES_IF:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool
100+
// CHECK-NEXT: %[[CONV_CAST_IF:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES_IF]] : !cir.bool to i1
101+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
102+
// CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i
103+
// CHECK-NEXT: %[[EQ_RES_SELF:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[TWO_LITERAL]]) : !s32i, !cir.bool
104+
// CHECK-NEXT: %[[CONV_CAST_SELF:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES_SELF]] : !cir.bool to i1
105+
// CHECK-NEXT: acc.kernels self(%[[CONV_CAST_SELF]]) if(%[[CONV_CAST_IF]]) {
106+
// CHECK-NEXT: acc.terminator
107+
// CHECK-NEXT: } loc
108+
67109
// CHECK-NEXT: cir.return
68110
}

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

Lines changed: 43 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
1+
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
22

33
void acc_parallel(int cond) {
44
// CHECK: cir.func @acc_parallel(%[[ARG:.*]]: !s32i{{.*}}) {
@@ -63,5 +63,47 @@ void acc_parallel(int cond) {
6363
// CHECK-NEXT: acc.yield
6464
// CHECK-NEXT: } loc
6565

66+
#pragma acc parallel if(cond)
67+
{}
68+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
69+
// CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
70+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
71+
// CHECK-NEXT: acc.parallel if(%[[CONV_CAST]]) {
72+
// CHECK-NEXT: acc.yield
73+
// CHECK-NEXT: } loc
74+
75+
#pragma acc parallel if(1)
76+
{}
77+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
78+
// CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ONE_LITERAL]] : !s32i), !cir.bool
79+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
80+
// CHECK-NEXT: acc.parallel if(%[[CONV_CAST]]) {
81+
// CHECK-NEXT: acc.yield
82+
// CHECK-NEXT: } loc
83+
84+
#pragma acc parallel if(cond == 1)
85+
{}
86+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
87+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
88+
// CHECK-NEXT: %[[EQ_RES:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool
89+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES]] : !cir.bool to i1
90+
// CHECK-NEXT: acc.parallel if(%[[CONV_CAST]]) {
91+
// CHECK-NEXT: acc.yield
92+
// CHECK-NEXT: } loc
93+
94+
#pragma acc parallel if(cond == 1) self(cond == 2)
95+
{}
96+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
97+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
98+
// CHECK-NEXT: %[[EQ_RES_IF:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool
99+
// CHECK-NEXT: %[[CONV_CAST_IF:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES_IF]] : !cir.bool to i1
100+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
101+
// CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i
102+
// CHECK-NEXT: %[[EQ_RES_SELF:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[TWO_LITERAL]]) : !s32i, !cir.bool
103+
// CHECK-NEXT: %[[CONV_CAST_SELF:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES_SELF]] : !cir.bool to i1
104+
// CHECK-NEXT: acc.parallel self(%[[CONV_CAST_SELF]]) if(%[[CONV_CAST_IF]]) {
105+
// CHECK-NEXT: acc.yield
106+
// CHECK-NEXT: } loc
107+
66108
// CHECK-NEXT: cir.return
67109
}

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

Lines changed: 43 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang_cc1 -fopenacc -emit-cir -fclangir %s -o - | FileCheck %s
1+
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
22

33
void acc_serial(int cond) {
44
// CHECK: cir.func @acc_serial(%[[ARG:.*]]: !s32i{{.*}}) {
@@ -64,5 +64,47 @@ void acc_serial(int cond) {
6464
// CHECK-NEXT: acc.yield
6565
// CHECK-NEXT: } loc
6666

67+
#pragma acc serial if(cond)
68+
{}
69+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
70+
// CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[COND_LOAD]] : !s32i), !cir.bool
71+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
72+
// CHECK-NEXT: acc.serial if(%[[CONV_CAST]]) {
73+
// CHECK-NEXT: acc.yield
74+
// CHECK-NEXT: } loc
75+
76+
#pragma acc serial if(1)
77+
{}
78+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
79+
// CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast(int_to_bool, %[[ONE_LITERAL]] : !s32i), !cir.bool
80+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
81+
// CHECK-NEXT: acc.serial if(%[[CONV_CAST]]) {
82+
// CHECK-NEXT: acc.yield
83+
// CHECK-NEXT: } loc
84+
85+
#pragma acc serial if(cond == 1)
86+
{}
87+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
88+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
89+
// CHECK-NEXT: %[[EQ_RES:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool
90+
// CHECK-NEXT: %[[CONV_CAST:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES]] : !cir.bool to i1
91+
// CHECK-NEXT: acc.serial if(%[[CONV_CAST]]) {
92+
// CHECK-NEXT: acc.yield
93+
// CHECK-NEXT: } loc
94+
95+
#pragma acc serial if(cond == 1) self(cond == 2)
96+
{}
97+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
98+
// CHECK-NEXT: %[[ONE_LITERAL:.*]] = cir.const #cir.int<1> : !s32i
99+
// CHECK-NEXT: %[[EQ_RES_IF:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[ONE_LITERAL]]) : !s32i, !cir.bool
100+
// CHECK-NEXT: %[[CONV_CAST_IF:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES_IF]] : !cir.bool to i1
101+
// CHECK-NEXT: %[[COND_LOAD:.*]] = cir.load %[[COND]] : !cir.ptr<!s32i>, !s32i
102+
// CHECK-NEXT: %[[TWO_LITERAL:.*]] = cir.const #cir.int<2> : !s32i
103+
// CHECK-NEXT: %[[EQ_RES_SELF:.*]] = cir.cmp(eq, %[[COND_LOAD]], %[[TWO_LITERAL]]) : !s32i, !cir.bool
104+
// CHECK-NEXT: %[[CONV_CAST_SELF:.*]] = builtin.unrealized_conversion_cast %[[EQ_RES_SELF]] : !cir.bool to i1
105+
// CHECK-NEXT: acc.serial self(%[[CONV_CAST_SELF]]) if(%[[CONV_CAST_IF]]) {
106+
// CHECK-NEXT: acc.yield
107+
// CHECK-NEXT: } loc
108+
67109
// CHECK-NEXT: cir.return
68110
}

0 commit comments

Comments
 (0)