Skip to content

Commit f7b4018

Browse files
erichkeaneerichkeane
authored
[OpenACC][CIR] Implement atomic update lowering (llvm#164836)
This is the 3rd of 4 forms of the 'atomic' construct. This one allows increment/decrement, compound-assign, and assign-to-bin-op(referencing the original variable). All of the above is enforced during Sema, but for our purposes, we ONLY need to know the variable on the LHS and the expression, so this does that. The ACC dialect for acc.atomic.update uses a 'recipe' as well, which takes the VALUE, and yields the value of the updated value. To simplify the implementation, our lowering very simply creates an alloca inside the recipe, stores the passed-in value, then loads/yields it at the end.
1 parent dddcb84 commit f7b4018

File tree

5 files changed

+253
-16
lines changed

5 files changed

+253
-16
lines changed

‎clang/lib/AST/StmtOpenACC.cpp‎

Lines changed: 36 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -326,35 +326,48 @@ OpenACCAtomicConstruct *OpenACCAtomicConstruct::Create(
326326

327327
static std::pair<const Expr *, const Expr *> getBinaryOpArgs(const Expr *Op) {
328328
if (const auto *BO = dyn_cast<BinaryOperator>(Op)) {
329-
assert(BO->getOpcode() == BO_Assign);
329+
assert(BO->isAssignmentOp());
330330
return {BO->getLHS(), BO->getRHS()};
331331
}
332332

333333
const auto *OO = cast<CXXOperatorCallExpr>(Op);
334-
assert(OO->getOperator() == OO_Equal);
335-
334+
assert(OO->isAssignmentOp());
336335
return {OO->getArg(0), OO->getArg(1)};
337336
}
338337

338+
static std::pair<bool, const Expr *> getUnaryOpArgs(const Expr *Op) {
339+
if (const auto *UO = dyn_cast<UnaryOperator>(Op))
340+
return {true, UO->getSubExpr()};
341+
342+
if (const auto *OpCall = dyn_cast<CXXOperatorCallExpr>(Op)) {
343+
// Post-inc/dec have a second unused argument to differentiate it, so we
344+
// accept -- or ++ as unary, or any operator call with only 1 arg.
345+
if (OpCall->getNumArgs() == 1 || OpCall->getOperator() != OO_PlusPlus ||
346+
OpCall->getOperator() != OO_MinusMinus)
347+
return {true, OpCall->getArg(0)};
348+
}
349+
350+
return {false, nullptr};
351+
}
352+
339353
const OpenACCAtomicConstruct::StmtInfo
340354
OpenACCAtomicConstruct::getAssociatedStmtInfo() const {
341355
// This ends up being a vastly simplified version of SemaOpenACCAtomic, since
342356
// it doesn't have to worry about erroring out, but we should do a lot of
343357
// asserts to ensure we don't get off into the weeds.
344358
assert(getAssociatedStmt() && "invalid associated stmt?");
345359

360+
const Expr *AssocStmt = cast<const Expr>(getAssociatedStmt());
346361
switch (AtomicKind) {
347-
case OpenACCAtomicKind::None:
348-
case OpenACCAtomicKind::Update:
349362
case OpenACCAtomicKind::Capture:
350-
assert(false && "Only 'read'/'write' have been implemented here");
363+
assert(false && "Only 'read'/'write'/'update' have been implemented here");
351364
return {};
352365
case OpenACCAtomicKind::Read: {
353366
// Read only supports the format 'v = x'; where both sides are a scalar
354367
// expression. This can come in 2 forms; BinaryOperator or
355368
// CXXOperatorCallExpr (rarely).
356369
std::pair<const Expr *, const Expr *> BinaryArgs =
357-
getBinaryOpArgs(cast<const Expr>(getAssociatedStmt()));
370+
getBinaryOpArgs(AssocStmt);
358371
// We want the L-value for each side, so we ignore implicit casts.
359372
return {BinaryArgs.first->IgnoreImpCasts(),
360373
BinaryArgs.second->IgnoreImpCasts(), /*expr=*/nullptr};
@@ -364,13 +377,28 @@ OpenACCAtomicConstruct::getAssociatedStmtInfo() const {
364377
// type, and 'x' is a scalar l value. As above, this can come in 2 forms;
365378
// Binary Operator or CXXOperatorCallExpr.
366379
std::pair<const Expr *, const Expr *> BinaryArgs =
367-
getBinaryOpArgs(cast<const Expr>(getAssociatedStmt()));
380+
getBinaryOpArgs(AssocStmt);
368381
// We want the L-value for ONLY the X side, so we ignore implicit casts. For
369382
// the right side (the expr), we emit it as an r-value so we need to
370383
// maintain implicit casts.
371384
return {/*v=*/nullptr, BinaryArgs.first->IgnoreImpCasts(),
372385
BinaryArgs.second};
373386
}
387+
case OpenACCAtomicKind::None:
388+
case OpenACCAtomicKind::Update: {
389+
std::pair<bool, const Expr *> UnaryArgs = getUnaryOpArgs(AssocStmt);
390+
if (UnaryArgs.first)
391+
return {/*v=*/nullptr, UnaryArgs.second->IgnoreImpCasts(),
392+
/*expr=*/nullptr};
393+
394+
std::pair<const Expr *, const Expr *> BinaryArgs =
395+
getBinaryOpArgs(AssocStmt);
396+
// For binary args, we just store the RHS as an expression (in the
397+
// expression slot), since the codegen just wants the whole thing for a
398+
// recipe.
399+
return {/*v=*/nullptr, BinaryArgs.first->IgnoreImpCasts(),
400+
BinaryArgs.second};
401+
}
374402
}
375403

376404
llvm_unreachable("unknown OpenACC atomic kind");

‎clang/lib/CIR/CodeGen/CIRGenFunction.h‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -665,6 +665,12 @@ class CIRGenFunction : public CIRGenTypeCache {
665665
symbolTable.insert(vd, addr.getPointer());
666666
}
667667

668+
// Replaces the address of the local variable, if it exists. Else does the
669+
// same thing as setAddrOfLocalVar.
670+
void replaceAddrOfLocalVar(const clang::VarDecl *vd, Address addr) {
671+
localDeclMap.insert_or_assign(vd, addr);
672+
}
673+
668674
// A class to allow reverting changes to a var-decl's registration to the
669675
// localDeclMap. This is used in cases where things are being inserted into
670676
// the variable list but don't follow normal lookup/search rules, like in

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

Lines changed: 58 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -304,12 +304,21 @@ CIRGenFunction::emitOpenACCCacheConstruct(const OpenACCCacheConstruct &s) {
304304
return mlir::success();
305305
}
306306

307+
const VarDecl *getLValueDecl(const Expr *e) {
308+
// We are going to assume that after stripping implicit casts, that the LValue
309+
// is just a DRE around the var-decl.
310+
311+
e = e->IgnoreImpCasts();
312+
313+
const auto *dre = cast<DeclRefExpr>(e);
314+
return cast<VarDecl>(dre->getDecl());
315+
}
316+
307317
mlir::LogicalResult
308318
CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {
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) {
319+
// For now, we are only support 'read'/'write'/'update', so diagnose. We can
320+
// switch on the kind later once we implement the 'capture' form.
321+
if (s.getAtomicKind() == OpenACCAtomicKind::Capture) {
313322
cgm.errorNYI(s.getSourceRange(), "OpenACC Atomic Construct");
314323
return mlir::failure();
315324
}
@@ -318,11 +327,10 @@ CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {
318327
// expression it is associated with rather than emitting it inside of it. So
319328
// it has custom emit logic.
320329
mlir::Location start = getLoc(s.getSourceRange().getBegin());
330+
mlir::Location end = getLoc(s.getSourceRange().getEnd());
321331
OpenACCAtomicConstruct::StmtInfo inf = s.getAssociatedStmtInfo();
322332

323333
switch (s.getAtomicKind()) {
324-
case OpenACCAtomicKind::None:
325-
case OpenACCAtomicKind::Update:
326334
case OpenACCAtomicKind::Capture:
327335
llvm_unreachable("Unimplemented atomic construct type, should have "
328336
"diagnosed/returned above");
@@ -353,6 +361,50 @@ CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {
353361
s.clauses());
354362
return mlir::success();
355363
}
364+
case OpenACCAtomicKind::None:
365+
case OpenACCAtomicKind::Update: {
366+
mlir::Value x = emitLValue(inf.X).getPointer();
367+
auto op =
368+
mlir::acc::AtomicUpdateOp::create(builder, start, x, /*ifCond=*/{});
369+
emitOpenACCClauses(op, s.getDirectiveKind(), s.getDirectiveLoc(),
370+
s.clauses());
371+
mlir::LogicalResult res = mlir::success();
372+
{
373+
mlir::OpBuilder::InsertionGuard guardCase(builder);
374+
mlir::Type argTy = cast<cir::PointerType>(x.getType()).getPointee();
375+
std::array<mlir::Type, 1> recipeType{argTy};
376+
std::array<mlir::Location, 1> recipeLoc{start};
377+
mlir::Block *recipeBlock = builder.createBlock(
378+
&op.getRegion(), op.getRegion().end(), recipeType, recipeLoc);
379+
builder.setInsertionPointToEnd(recipeBlock);
380+
381+
// Since we have an initial value that we know is a scalar type, we can
382+
// just emit the entire statement here after sneaking-in our 'alloca' in
383+
// the right place, then loading out of it. Flang does a lot less work
384+
// (probably does its own emitting!), but we have more complicated AST
385+
// nodes to worry about, so we can just count on opt to remove the extra
386+
// alloca/load/store set.
387+
auto alloca = cir::AllocaOp::create(
388+
builder, start, x.getType(), argTy, "x_var",
389+
cgm.getSize(getContext().getTypeAlignInChars(inf.X->getType())));
390+
391+
alloca.setInitAttr(mlir::UnitAttr::get(&getMLIRContext()));
392+
builder.CIRBaseBuilderTy::createStore(start, recipeBlock->getArgument(0),
393+
alloca);
394+
395+
const VarDecl *xval = getLValueDecl(inf.X);
396+
CIRGenFunction::DeclMapRevertingRAII declMapRAII{*this, xval};
397+
replaceAddrOfLocalVar(
398+
xval, Address{alloca, argTy, getContext().getDeclAlign(xval)});
399+
400+
res = emitStmt(s.getAssociatedStmt(), /*useCurrentScope=*/true);
401+
402+
auto load = cir::LoadOp::create(builder, start, {alloca});
403+
mlir::acc::YieldOp::create(builder, end, {load});
404+
}
405+
406+
return res;
407+
}
356408
}
357409

358410
llvm_unreachable("unknown OpenACC atomic kind");

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

Lines changed: 151 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,151 @@
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+
struct HasOps {
4+
operator float();
5+
int thing();
6+
};
7+
8+
void use(int x, unsigned int y, float f, HasOps ops) {
9+
// CHECK: cir.func{{.*}}(%[[X_ARG:.*]]: !s32i{{.*}}, %[[Y_ARG:.*]]: !u32i{{.*}}, %[[F_ARG:.*]]: !cir.float{{.*}}){{.*}}, %[[OPS_ARG:.*]]: !rec_HasOps{{.*}}) {
10+
// CHECK-NEXT: %[[X_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init]
11+
// CHECK-NEXT: %[[Y_ALLOCA:.*]] = cir.alloca !u32i, !cir.ptr<!u32i>, ["y", init]
12+
// CHECK-NEXT: %[[F_ALLOCA:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["f", init]
13+
// CHECK-NEXT: %[[OPS_ALLOCA:.*]] = cir.alloca !rec_HasOps, !cir.ptr<!rec_HasOps>, ["ops", init]
14+
// CHECK-NEXT: cir.store %[[X_ARG]], %[[X_ALLOCA]] : !s32i, !cir.ptr<!s32i>
15+
// CHECK-NEXT: cir.store %[[Y_ARG]], %[[Y_ALLOCA]] : !u32i, !cir.ptr<!u32i>
16+
// CHECK-NEXT: cir.store %[[F_ARG]], %[[F_ALLOCA]] : !cir.float, !cir.ptr<!cir.float>
17+
// CHECK-NEXT: cir.store %[[OPS_ARG]], %[[OPS_ALLOCA]] : !rec_HasOps, !cir.ptr<!rec_HasOps>
18+
19+
// CHECK-NEXT: acc.atomic.update %[[X_ALLOCA]] : !cir.ptr<!s32i> {
20+
// CHECK-NEXT: ^bb0(%[[RECIPE_ARG:.*]]: !s32i{{.*}}):
21+
// CHECK-NEXT: %[[TEMP_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x_var", init]
22+
// CHECK-NEXT: cir.store %[[RECIPE_ARG]], %[[TEMP_ALLOCA]] : !s32i, !cir.ptr<!s32i>
23+
//
24+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!s32i>, !s32i
25+
// CHECK-NEXT: %[[INC:.*]] = cir.unary(inc, %[[TEMP_LOAD]]) nsw : !s32i, !s32i
26+
// CHECK-NEXT: cir.store {{.*}}%[[INC]], %[[TEMP_ALLOCA]] : !s32i, !cir.ptr<!s32i>
27+
//
28+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!s32i>, !s32i
29+
// CHECK-NEXT: acc.yield %[[TEMP_LOAD]] : !s32i
30+
// CHECK-NEXT: }
31+
#pragma acc atomic update
32+
++x;
33+
34+
// CHECK-NEXT: acc.atomic.update %[[Y_ALLOCA]] : !cir.ptr<!u32i> {
35+
// CHECK-NEXT: ^bb0(%[[RECIPE_ARG:.*]]: !u32i{{.*}}):
36+
// CHECK-NEXT: %[[TEMP_ALLOCA:.*]] = cir.alloca !u32i, !cir.ptr<!u32i>, ["x_var", init]
37+
// CHECK-NEXT: cir.store %[[RECIPE_ARG]], %[[TEMP_ALLOCA]] : !u32i, !cir.ptr<!u32i>
38+
//
39+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!u32i>, !u32i
40+
// CHECK-NEXT: %[[INC:.*]] = cir.unary(inc, %[[TEMP_LOAD]]) : !u32i, !u32i
41+
// CHECK-NEXT: cir.store {{.*}}%[[INC]], %[[TEMP_ALLOCA]] : !u32i, !cir.ptr<!u32i>
42+
//
43+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!u32i>, !u32i
44+
// CHECK-NEXT: acc.yield %[[TEMP_LOAD]] : !u32i
45+
// CHECK-NEXT: }
46+
#pragma acc atomic update
47+
y++;
48+
49+
// CHECK-NEXT: acc.atomic.update %[[F_ALLOCA]] : !cir.ptr<!cir.float> {
50+
// CHECK-NEXT: ^bb0(%[[RECIPE_ARG:.*]]: !cir.float{{.*}}):
51+
// CHECK-NEXT: %[[TEMP_ALLOCA:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["x_var", init]
52+
// CHECK-NEXT: cir.store %[[RECIPE_ARG]], %[[TEMP_ALLOCA]] : !cir.float, !cir.ptr<!cir.float>
53+
//
54+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!cir.float>, !cir.float
55+
// CHECK-NEXT: %[[INC:.*]] = cir.unary(dec, %[[TEMP_LOAD]]) : !cir.float, !cir.float
56+
// CHECK-NEXT: cir.store {{.*}}%[[INC]], %[[TEMP_ALLOCA]] : !cir.float, !cir.ptr<!cir.float>
57+
//
58+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!cir.float>, !cir.float
59+
// CHECK-NEXT: acc.yield %[[TEMP_LOAD]] : !cir.float
60+
// CHECK-NEXT: }
61+
#pragma acc atomic update
62+
f--;
63+
64+
// CHECK-NEXT: acc.atomic.update %[[X_ALLOCA]] : !cir.ptr<!s32i> {
65+
// CHECK-NEXT: ^bb0(%[[RECIPE_ARG:.*]]: !s32i{{.*}}):
66+
// CHECK-NEXT: %[[TEMP_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x_var", init]
67+
// CHECK-NEXT: cir.store %[[RECIPE_ARG]], %[[TEMP_ALLOCA]] : !s32i, !cir.ptr<!s32i>
68+
//
69+
// CHECK-NEXT: %[[F_LOAD:.*]] = cir.load{{.*}} %[[F_ALLOCA]] : !cir.ptr<!cir.float>, !cir.float
70+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!s32i>, !s32i
71+
// CHECK-NEXT: %[[INT_TO_F:.*]] = cir.cast int_to_float %[[TEMP_LOAD]] : !s32i -> !cir.float
72+
// CHECK-NEXT: %[[ADD:.*]] = cir.binop(add, %[[INT_TO_F]], %[[F_LOAD]]) : !cir.float
73+
// CHECK-NEXT: %[[F_TO_INT:.*]] = cir.cast float_to_int %[[ADD]] : !cir.float -> !s32i
74+
// CHECK-NEXT: cir.store{{.*}} %[[F_TO_INT]], %[[TEMP_ALLOCA]] : !s32i, !cir.ptr<!s32i>
75+
//
76+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!s32i>, !s32i
77+
// CHECK-NEXT: acc.yield %[[TEMP_LOAD]] : !s32i
78+
// CHECK-NEXT: }
79+
#pragma acc atomic update
80+
x += f;
81+
82+
// CHECK-NEXT: acc.atomic.update %[[F_ALLOCA]] : !cir.ptr<!cir.float> {
83+
// CHECK-NEXT: ^bb0(%[[RECIPE_ARG:.*]]: !cir.float{{.*}}):
84+
// CHECK-NEXT: %[[TEMP_ALLOCA:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["x_var", init]
85+
// CHECK-NEXT: cir.store %[[RECIPE_ARG]], %[[TEMP_ALLOCA]] : !cir.float, !cir.ptr<!cir.float>
86+
//
87+
// CHECK-NEXT: %[[Y_LOAD:.*]] = cir.load{{.*}} %[[Y_ALLOCA]] : !cir.ptr<!u32i>, !u32i
88+
// CHECK-NEXT: %[[INT_TO_F:.*]] = cir.cast int_to_float %[[Y_LOAD]] : !u32i -> !cir.float
89+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!cir.float>, !cir.float
90+
// CHECK-NEXT: %[[DIV:.*]] = cir.binop(div, %[[TEMP_LOAD]], %[[INT_TO_F]]) : !cir.float
91+
// CHECK-NEXT: cir.store{{.*}} %[[DIV]], %[[TEMP_ALLOCA]] : !cir.float, !cir.ptr<!cir.float>
92+
//
93+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!cir.float>, !cir.float
94+
// CHECK-NEXT: acc.yield %[[TEMP_LOAD]] : !cir.float
95+
// CHECK-NEXT: }
96+
#pragma acc atomic update
97+
f /= y;
98+
99+
// CHECK-NEXT: acc.atomic.update %[[Y_ALLOCA]] : !cir.ptr<!u32i> {
100+
// CHECK-NEXT: ^bb0(%[[RECIPE_ARG:.*]]: !u32i{{.*}}):
101+
// CHECK-NEXT: %[[TEMP_ALLOCA:.*]] = cir.alloca !u32i, !cir.ptr<!u32i>, ["x_var", init]
102+
// CHECK-NEXT: cir.store %[[RECIPE_ARG]], %[[TEMP_ALLOCA]] : !u32i, !cir.ptr<!u32i>
103+
//
104+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!u32i>, !u32i
105+
// CHECK-NEXT: %[[CALL:.*]] = cir.call {{.*}}(%[[OPS_ALLOCA]]) : (!cir.ptr<!rec_HasOps>) -> !s32i
106+
// CHECK-NEXT: %[[CALL_CAST:.*]] = cir.cast integral %[[CALL]] : !s32i -> !u32i
107+
// CHECK-NEXT: %[[MUL:.*]] = cir.binop(mul, %[[TEMP_LOAD]], %[[CALL_CAST]]) : !u32i
108+
// CHECK-NEXT: cir.store{{.*}} %[[MUL]], %[[TEMP_ALLOCA]] : !u32i, !cir.ptr<!u32i>
109+
//
110+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!u32i>, !u32i
111+
// CHECK-NEXT: acc.yield %[[TEMP_LOAD]] : !u32i
112+
// CHECK-NEXT: }
113+
114+
#pragma acc atomic update
115+
y = y * ops.thing();
116+
117+
// CHECK-NEXT: acc.atomic.update %[[X_ALLOCA]] : !cir.ptr<!s32i> {
118+
// CHECK-NEXT: ^bb0(%[[RECIPE_ARG:.*]]: !s32i{{.*}}):
119+
// CHECK-NEXT: %[[TEMP_ALLOCA:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x_var", init]
120+
// CHECK-NEXT: cir.store %[[RECIPE_ARG]], %[[TEMP_ALLOCA]] : !s32i, !cir.ptr<!s32i>
121+
//
122+
// CHECK-NEXT: %[[CALL:.*]] = cir.call {{.*}}(%[[OPS_ALLOCA]]) : (!cir.ptr<!rec_HasOps>) -> !s32i
123+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!s32i>, !s32i
124+
// CHECK-NEXT: %[[OR:.*]] = cir.binop(or, %[[CALL]], %[[INT_TO_F]]) : !s32i
125+
// CHECK-NEXT: cir.store{{.*}} %[[OR]], %[[TEMP_ALLOCA]] : !s32i, !cir.ptr<!s32i>
126+
//
127+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!s32i>, !s32i
128+
// CHECK-NEXT: acc.yield %[[TEMP_LOAD]] : !s32i
129+
// CHECK-NEXT: }
130+
#pragma acc atomic update
131+
x = ops.thing() | x;
132+
133+
// CHECK-NEXT: %[[X_LOAD:.*]] = cir.load{{.*}} %[[X_ALLOCA]] : !cir.ptr<!s32i>, !s32i
134+
// CHECK-NEXT: %[[BOOL_CAST:.*]] = cir.cast int_to_bool %[[X_LOAD]] : !s32i -> !cir.bool
135+
// CHECK-NEXT: %[[X_CAST:.*]] = builtin.unrealized_conversion_cast %[[BOOL_CAST]] : !cir.bool to i1
136+
// CHECK-NEXT: acc.atomic.update if(%[[X_CAST]]) %[[F_ALLOCA]] : !cir.ptr<!cir.float> {
137+
// CHECK-NEXT: ^bb0(%[[RECIPE_ARG:.*]]: !cir.float{{.*}}):
138+
// CHECK-NEXT: %[[TEMP_ALLOCA:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["x_var", init]
139+
// CHECK-NEXT: cir.store %[[RECIPE_ARG]], %[[TEMP_ALLOCA]] : !cir.float, !cir.ptr<!cir.float>
140+
//
141+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!cir.float>, !cir.float
142+
// CHECK-NEXT: %[[CALL:.*]] = cir.call {{.*}}(%[[OPS_ALLOCA]]) : (!cir.ptr<!rec_HasOps>) -> !cir.float
143+
// CHECK-NEXT: %[[SUB:.*]] = cir.binop(sub, %[[TEMP_LOAD]], %[[CALL]]) : !cir.float
144+
// CHECK-NEXT: cir.store{{.*}} %[[SUB]], %[[TEMP_ALLOCA]] : !cir.float, !cir.ptr<!cir.float>
145+
//
146+
// CHECK-NEXT: %[[TEMP_LOAD:.*]] = cir.load{{.*}} %[[TEMP_ALLOCA]] : !cir.ptr<!cir.float>, !cir.float
147+
// CHECK-NEXT: acc.yield %[[TEMP_LOAD]] : !cir.float
148+
// CHECK-NEXT: }
149+
#pragma acc atomic update if (x)
150+
f = f - ops;
151+
}

‎clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,8 +3,8 @@
33
void HelloWorld(int *A, int *B, int *C, int N) {
44

55
// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Atomic Construct}}
6-
#pragma acc atomic
7-
N = N + 1;
6+
#pragma acc atomic capture
7+
B = A += ++N;
88

99
// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
1010
#pragma acc declare create(A)

0 commit comments

Comments
 (0)