Skip to content

Commit c0fa2b6

Browse files
erichkeaneLukacma
erichkeane
authored and
Lukacma
committed
[OpenACC][CIR] Implement atomic-write lowering (llvm#164627)
This is a slightly more complicated variant of this, which supports 'x = expr', so the right hand side is an r-value. This patch implements that, adds some tests, and does some minor refactoring to the infrastructure added for the 'atomic read' to make it more flexible for 'write'. This is the second of four 'atomic' kinds.
1 parent 31a9573 commit c0fa2b6

File tree

4 files changed

+129
-28
lines changed

4 files changed

+129
-28
lines changed

‎clang/include/clang/AST/StmtOpenACC.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -821,6 +821,7 @@ class OpenACCAtomicConstruct final
821821
struct StmtInfo {
822822
const Expr *V;
823823
const Expr *X;
824+
const Expr *Expr;
824825
// TODO: OpenACC: We should expand this as we're implementing the other
825826
// atomic construct kinds.
826827
};

‎clang/lib/AST/StmtOpenACC.cpp‎

Lines changed: 32 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -324,6 +324,18 @@ OpenACCAtomicConstruct *OpenACCAtomicConstruct::Create(
324324
return Inst;
325325
}
326326

327+
static std::pair<const Expr *, const Expr *> getBinaryOpArgs(const Expr *Op) {
328+
if (const auto *BO = dyn_cast<BinaryOperator>(Op)) {
329+
assert(BO->getOpcode() == BO_Assign);
330+
return {BO->getLHS(), BO->getRHS()};
331+
}
332+
333+
const auto *OO = cast<CXXOperatorCallExpr>(Op);
334+
assert(OO->getOperator() == OO_Equal);
335+
336+
return {OO->getArg(0), OO->getArg(1)};
337+
}
338+
327339
const OpenACCAtomicConstruct::StmtInfo
328340
OpenACCAtomicConstruct::getAssociatedStmtInfo() const {
329341
// This ends up being a vastly simplified version of SemaOpenACCAtomic, since
@@ -333,27 +345,35 @@ OpenACCAtomicConstruct::getAssociatedStmtInfo() const {
333345

334346
switch (AtomicKind) {
335347
case OpenACCAtomicKind::None:
336-
case OpenACCAtomicKind::Write:
337348
case OpenACCAtomicKind::Update:
338349
case OpenACCAtomicKind::Capture:
339-
assert(false && "Only 'read' has been implemented here");
350+
assert(false && "Only 'read'/'write' have been implemented here");
340351
return {};
341352
case OpenACCAtomicKind::Read: {
342353
// Read only supports the format 'v = x'; where both sides are a scalar
343354
// expression. This can come in 2 forms; BinaryOperator or
344355
// CXXOperatorCallExpr (rarely).
345-
const Expr *AssignExpr = cast<const Expr>(getAssociatedStmt());
346-
if (const auto *BO = dyn_cast<BinaryOperator>(AssignExpr)) {
347-
assert(BO->getOpcode() == BO_Assign);
348-
return {BO->getLHS()->IgnoreImpCasts(), BO->getRHS()->IgnoreImpCasts()};
349-
}
350-
351-
const auto *OO = cast<CXXOperatorCallExpr>(AssignExpr);
352-
assert(OO->getOperator() == OO_Equal);
353-
354-
return {OO->getArg(0)->IgnoreImpCasts(), OO->getArg(1)->IgnoreImpCasts()};
356+
std::pair<const Expr *, const Expr *> BinaryArgs =
357+
getBinaryOpArgs(cast<const Expr>(getAssociatedStmt()));
358+
// We want the L-value for each side, so we ignore implicit casts.
359+
return {BinaryArgs.first->IgnoreImpCasts(),
360+
BinaryArgs.second->IgnoreImpCasts(), /*expr=*/nullptr};
355361
}
362+
case OpenACCAtomicKind::Write: {
363+
// Write supports only the format 'x = expr', where the expression is scalar
364+
// type, and 'x' is a scalar l value. As above, this can come in 2 forms;
365+
// Binary Operator or CXXOperatorCallExpr.
366+
std::pair<const Expr *, const Expr *> BinaryArgs =
367+
getBinaryOpArgs(cast<const Expr>(getAssociatedStmt()));
368+
// We want the L-value for ONLY the X side, so we ignore implicit casts. For
369+
// the right side (the expr), we emit it as an r-value so we need to
370+
// maintain implicit casts.
371+
return {/*v=*/nullptr, BinaryArgs.first->IgnoreImpCasts(),
372+
BinaryArgs.second};
356373
}
374+
}
375+
376+
llvm_unreachable("unknown OpenACC atomic kind");
357377
}
358378

359379
OpenACCCacheConstruct *OpenACCCacheConstruct::CreateEmpty(const ASTContext &C,

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

Lines changed: 41 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -306,9 +306,10 @@ CIRGenFunction::emitOpenACCCacheConstruct(const OpenACCCacheConstruct &s) {
306306

307307
mlir::LogicalResult
308308
CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {
309-
// For now, we are only support 'read', so diagnose. We can switch on the kind
310-
// later once we start implementing the other 3 forms.
311-
if (s.getAtomicKind() != OpenACCAtomicKind::Read) {
309+
// For now, we are only support 'read'/'write', so diagnose. We can switch on
310+
// the kind later once we start implementing the other 2 forms. While we
311+
if (s.getAtomicKind() != OpenACCAtomicKind::Read &&
312+
s.getAtomicKind() != OpenACCAtomicKind::Write) {
312313
cgm.errorNYI(s.getSourceRange(), "OpenACC Atomic Construct");
313314
return mlir::failure();
314315
}
@@ -318,17 +319,41 @@ CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {
318319
// it has custom emit logic.
319320
mlir::Location start = getLoc(s.getSourceRange().getBegin());
320321
OpenACCAtomicConstruct::StmtInfo inf = s.getAssociatedStmtInfo();
321-
// Atomic 'read' only permits 'v = x', where v and x are both scalar L values.
322-
// The getAssociatedStmtInfo strips off implicit casts, which includes
323-
// implicit conversions and L-to-R-Value conversions, so we can just emit it
324-
// as an L value. The Flang implementation has no problem with different
325-
// types, so it appears that the dialect can handle the conversions.
326-
mlir::Value v = emitLValue(inf.V).getPointer();
327-
mlir::Value x = emitLValue(inf.X).getPointer();
328-
mlir::Type resTy = convertType(inf.V->getType());
329-
auto op = mlir::acc::AtomicReadOp::create(builder, start, x, v, resTy,
330-
/*ifCond=*/{});
331-
emitOpenACCClauses(op, s.getDirectiveKind(), s.getDirectiveLoc(),
332-
s.clauses());
333-
return mlir::success();
322+
323+
switch (s.getAtomicKind()) {
324+
case OpenACCAtomicKind::None:
325+
case OpenACCAtomicKind::Update:
326+
case OpenACCAtomicKind::Capture:
327+
llvm_unreachable("Unimplemented atomic construct type, should have "
328+
"diagnosed/returned above");
329+
return mlir::failure();
330+
case OpenACCAtomicKind::Read: {
331+
332+
// Atomic 'read' only permits 'v = x', where v and x are both scalar L
333+
// values. The getAssociatedStmtInfo strips off implicit casts, which
334+
// includes implicit conversions and L-to-R-Value conversions, so we can
335+
// just emit it as an L value. The Flang implementation has no problem with
336+
// different types, so it appears that the dialect can handle the
337+
// conversions.
338+
mlir::Value v = emitLValue(inf.V).getPointer();
339+
mlir::Value x = emitLValue(inf.X).getPointer();
340+
mlir::Type resTy = convertType(inf.V->getType());
341+
auto op = mlir::acc::AtomicReadOp::create(builder, start, x, v, resTy,
342+
/*ifCond=*/{});
343+
emitOpenACCClauses(op, s.getDirectiveKind(), s.getDirectiveLoc(),
344+
s.clauses());
345+
return mlir::success();
346+
}
347+
case OpenACCAtomicKind::Write: {
348+
mlir::Value x = emitLValue(inf.X).getPointer();
349+
mlir::Value expr = emitAnyExpr(inf.Expr).getValue();
350+
auto op = mlir::acc::AtomicWriteOp::create(builder, start, x, expr,
351+
/*ifCond=*/{});
352+
emitOpenACCClauses(op, s.getDirectiveKind(), s.getDirectiveLoc(),
353+
s.clauses());
354+
return mlir::success();
355+
}
356+
}
357+
358+
llvm_unreachable("unknown OpenACC atomic kind");
334359
}

‎clang/test/CIR/CodeGenOpenACC/atomic-write.cpp‎

Lines changed: 55 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,55 @@
1+
// RUN: %clang_cc1 -fopenacc -triple x86_64-linux-gnu -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir -triple x86_64-linux-pc %s -o - | FileCheck %s
2+
3+
extern "C" bool condition(int x, unsigned int y, float f);
4+
extern "C" double do_thing(float f);
5+
6+
struct ConvertsToScalar {
7+
operator float();
8+
};
9+
10+
void use(int x, unsigned int y, float f, ConvertsToScalar cts) {
11+
// CHECK: cir.func{{.*}}(%[[X_ARG:.*]]: !s32i{{.*}}, %[[Y_ARG:.*]]: !u32i{{.*}}, %[[F_ARG:.*]]: !cir.float{{.*}}){{.*}}, %[[CTS_ARG:.*]]: !rec_ConvertsToScalar{{.*}}) {
12+
// CHECK-NEXT: %[[X_ALLOC:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init]
13+
// CHECK-NEXT: %[[Y_ALLOC:.*]] = cir.alloca !u32i, !cir.ptr<!u32i>, ["y", init]
14+
// CHECK-NEXT: %[[F_ALLOC:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["f", init]
15+
// CHECK-NEXT: %[[CTS_ALLOC:.*]] = cir.alloca !rec_ConvertsToScalar, !cir.ptr<!rec_ConvertsToScalar>, ["cts", init]
16+
//
17+
// CHECK-NEXT: cir.store %[[X_ARG]], %[[X_ALLOC]] : !s32i, !cir.ptr<!s32i>
18+
// CHECK-NEXT: cir.store %[[Y_ARG]], %[[Y_ALLOC]] : !u32i, !cir.ptr<!u32i>
19+
// CHECK-NEXT: cir.store %[[F_ARG]], %[[F_ALLOC]] : !cir.float, !cir.ptr<!cir.float>
20+
// CHECK-NEXT: cir.store %[[CTS_ARG]], %[[CTS_ALLOC]] : !rec_ConvertsToScalar, !cir.ptr<!rec_ConvertsToScalar>
21+
22+
// CHECK-NEXT: %[[Y_LOAD:.*]] = cir.load {{.*}}%[[Y_ALLOC]] : !cir.ptr<!u32i>, !u32i
23+
// CHECK-NEXT: %[[Y_TO_FLOAT:.*]] = cir.cast int_to_float %[[Y_LOAD]] : !u32i -> !cir.float
24+
// CHECK-NEXT: %[[F_LOAD:.*]] = cir.load {{.*}}%[[F_ALLOC]] : !cir.ptr<!cir.float>, !cir.float
25+
// CHECK-NEXT: %[[MUL:.*]] = cir.binop(mul, %[[Y_TO_FLOAT]], %[[F_LOAD]]) : !cir.float
26+
// CHECK-NEXT: %[[RHS_CAST:.*]] = cir.cast float_to_int %[[MUL]] : !cir.float -> !s32i
27+
// CHECK-NEXT: acc.atomic.write %[[X_ALLOC]] = %[[RHS_CAST]] : !cir.ptr<!s32i>, !s32i
28+
#pragma acc atomic write
29+
x = y * f;
30+
31+
// CHECK-NEXT: %[[F_LOAD:.*]] = cir.load {{.*}}%[[F_ALLOC]] : !cir.ptr<!cir.float>, !cir.float
32+
// CHECK-NEXT: %[[CALL:.*]] = cir.call @do_thing(%[[F_LOAD]]) : (!cir.float) -> !cir.double
33+
// CHECK-NEXT: %[[CALL_CAST:.*]] = cir.cast float_to_int %[[CALL]] : !cir.double -> !u32i
34+
// CHECK-NEXT: acc.atomic.write %[[Y_ALLOC]] = %[[CALL_CAST]] : !cir.ptr<!u32i>, !u32i
35+
#pragma acc atomic write
36+
y = do_thing(f);
37+
38+
// CHECK-NEXT: %[[X_LOAD:.*]] = cir.load {{.*}}%[[X_ALLOC]] : !cir.ptr<!s32i>, !s32i
39+
// CHECK-NEXT: %[[X_CAST:.*]] = cir.cast int_to_float %[[X_LOAD]] : !s32i -> !cir.float
40+
// CHECK-NEXT: %[[THING_CALL:.*]] = cir.call @do_thing(%[[X_CAST]]) : (!cir.float) -> !cir.double
41+
// CHECK-NEXT: %[[THING_CAST:.*]] = cir.cast floating %[[THING_CALL]] : !cir.double -> !cir.float
42+
// CHECK-NEXT: %[[X_LOAD:.*]] = cir.load {{.*}}%[[X_ALLOC]] : !cir.ptr<!s32i>, !s32i
43+
// CHECK-NEXT: %[[Y_LOAD:.*]] = cir.load {{.*}}%[[Y_ALLOC]] : !cir.ptr<!u32i>, !u32i
44+
// CHECK-NEXT: %[[F_LOAD:.*]] = cir.load {{.*}}%[[F_ALLOC]] : !cir.ptr<!cir.float>, !cir.float
45+
// CHECK-NEXT: %[[COND_CALL:.*]] = cir.call @condition(%[[X_LOAD]], %[[Y_LOAD]], %[[F_LOAD]]) : (!s32i, !u32i, !cir.float) -> !cir.bool
46+
// CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_CALL]] : !cir.bool to i1
47+
// CHECK-NEXT: acc.atomic.write if(%[[COND_CAST]]) %[[F_ALLOC]] = %[[THING_CAST]] : !cir.ptr<!cir.float>, !cir.float
48+
#pragma acc atomic write if (condition(x, y, f))
49+
f = do_thing(x);
50+
51+
// CHECK-NEXT: %[[CTS_CONV_CALL:.*]] = cir.call @{{.*}}(%[[CTS_ALLOC]]) : (!cir.ptr<!rec_ConvertsToScalar>) -> !cir.float
52+
// CHECK-NEXT: acc.atomic.write %[[F_ALLOC]] = %[[CTS_CONV_CALL]] : !cir.ptr<!cir.float>, !cir.float
53+
#pragma acc atomic write
54+
f = cts;
55+
}

0 commit comments

Comments
 (0)