[clang] [OpenACC][CIR] Fix atomic-capture single-line-postfix (PR #168717)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Nov 19 06:57:01 PST 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clangir
Author: Erich Keane (erichkeane)
<details>
<summary>Changes</summary>
In my last patch, it became clear during code review that the postfix operation was actually a read THEN update, not update/read like other single line versions. It wasn't clear at the time how much additional work this would be to make postfix work correctly (and they are a bit of a 'special' thing in codegen anyway), so this patch adds some functionality to sense this and special-cases it when generating the statement info for capture.
---
Full diff: https://github.com/llvm/llvm-project/pull/168717.diff
3 Files Affected:
- (modified) clang/include/clang/AST/StmtOpenACC.h (+9-3)
- (modified) clang/lib/AST/StmtOpenACC.cpp (+18-17)
- (modified) clang/test/CIR/CodeGenOpenACC/atomic-capture.cpp (+2-2)
``````````diff
diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h
index ad4e2d65771b8..2bd0b52071697 100644
--- a/clang/include/clang/AST/StmtOpenACC.h
+++ b/clang/include/clang/AST/StmtOpenACC.h
@@ -829,8 +829,13 @@ class OpenACCAtomicConstruct final
// Listed as 'expr' in the standard, this is typically a generic expression
// as a component.
const Expr *RefExpr;
+ // If this is an 'update', records whether this is a post-fix
+ // increment/decrement. In the case where we have a single-line variant of
+ // 'capture' we have to form the IR differently if this is the case to make
+ // sure the old value is 'read' in the 2nd step.
+ bool IsPostfixIncDec = false;
static SingleStmtInfo Empty() {
- return {nullptr, nullptr, nullptr, nullptr};
+ return {nullptr, nullptr, nullptr, nullptr, false};
}
static SingleStmtInfo createRead(const Expr *WholeExpr, const Expr *V,
@@ -841,8 +846,9 @@ class OpenACCAtomicConstruct final
const Expr *RefExpr) {
return {WholeExpr, /*V=*/nullptr, X, RefExpr};
}
- static SingleStmtInfo createUpdate(const Expr *WholeExpr, const Expr *X) {
- return {WholeExpr, /*V=*/nullptr, X, /*RefExpr=*/nullptr};
+ static SingleStmtInfo createUpdate(const Expr *WholeExpr, const Expr *X,
+ bool PostfixIncDec) {
+ return {WholeExpr, /*V=*/nullptr, X, /*RefExpr=*/nullptr, PostfixIncDec};
}
};
diff --git a/clang/lib/AST/StmtOpenACC.cpp b/clang/lib/AST/StmtOpenACC.cpp
index d3a7e7601f618..ec8ceb949c6c0 100644
--- a/clang/lib/AST/StmtOpenACC.cpp
+++ b/clang/lib/AST/StmtOpenACC.cpp
@@ -347,16 +347,17 @@ getBinaryAssignOpArgs(const Expr *Op) {
return getBinaryAssignOpArgs(Op, IsCompoundAssign);
}
-static std::optional<const Expr *> getUnaryOpArgs(const Expr *Op) {
+static std::optional<std::pair<const Expr *, bool>>
+getUnaryOpArgs(const Expr *Op) {
if (const auto *UO = dyn_cast<UnaryOperator>(Op))
- return UO->getSubExpr();
+ return {{UO->getSubExpr(), UO->isPostfix()}};
if (const auto *OpCall = dyn_cast<CXXOperatorCallExpr>(Op)) {
// Post-inc/dec have a second unused argument to differentiate it, so we
// accept -- or ++ as unary, or any operator call with only 1 arg.
if (OpCall->getNumArgs() == 1 || OpCall->getOperator() == OO_PlusPlus ||
OpCall->getOperator() == OO_MinusMinus)
- return {OpCall->getArg(0)};
+ return {{OpCall->getArg(0), /*IsPostfix=*/OpCall->getNumArgs() == 1}};
}
return std::nullopt;
@@ -410,10 +411,10 @@ getWriteStmtInfo(const Expr *E) {
static std::optional<OpenACCAtomicConstruct::SingleStmtInfo>
getUpdateStmtInfo(const Expr *E) {
- std::optional<const Expr *> UnaryArgs = getUnaryOpArgs(E);
+ std::optional<std::pair<const Expr *, bool>> UnaryArgs = getUnaryOpArgs(E);
if (UnaryArgs) {
auto Res = OpenACCAtomicConstruct::SingleStmtInfo::createUpdate(
- E, (*UnaryArgs)->IgnoreImpCasts());
+ E, UnaryArgs->first->IgnoreImpCasts(), UnaryArgs->second);
if (!Res.X->isLValue() || !Res.X->getType()->isScalarType())
return std::nullopt;
@@ -428,7 +429,7 @@ getUpdateStmtInfo(const Expr *E) {
return std::nullopt;
auto Res = OpenACCAtomicConstruct::SingleStmtInfo::createUpdate(
- E, BinaryArgs->first->IgnoreImpCasts());
+ E, BinaryArgs->first->IgnoreImpCasts(), /*PostFixIncDec=*/false);
if (!Res.X->isLValue() || !Res.X->getType()->isScalarType())
return std::nullopt;
@@ -513,17 +514,12 @@ getCaptureStmtInfo(const Stmt *AssocStmt) {
return OpenACCAtomicConstruct::StmtInfo::createUpdateRead(*Update, *Read);
} else {
- // All of the possible forms (listed below) that are writable as a single
- // line are expressed as an update, then as a read. We should be able to
- // just run these two in the right order.
- // UPDATE: READ
- // v = x++;
- // v = x--;
- // v = ++x;
- // v = --x;
- // v = x binop=expr
- // v = x = x binop expr
- // v = x = expr binop x
+ // All of the forms that can be done in a single line fall into 2
+ // categories: update/read, or read/update. The special cases are the
+ // postfix unary operators, which we have to make sure we do the 'read'
+ // first. However, we still parse these as the RHS first, so we have a
+ // 'reversing' step. READ: UPDATE v = x++; v = x--; UPDATE: READ v = ++x; v
+ // = --x; v = x binop=expr v = x = x binop expr v = x = expr binop x
const Expr *E = cast<const Expr>(AssocStmt);
@@ -535,6 +531,11 @@ getCaptureStmtInfo(const Stmt *AssocStmt) {
// Fixup this, since the 'X' for the read is the result after write, but is
// the same value as the LHS-most variable of the update(its X).
Read->X = Update->X;
+
+ // Postfix is a read FIRST, then an update.
+ if (Update->IsPostfixIncDec)
+ return OpenACCAtomicConstruct::StmtInfo::createReadUpdate(*Read, *Update);
+
return OpenACCAtomicConstruct::StmtInfo::createUpdateRead(*Update, *Read);
}
return {};
diff --git a/clang/test/CIR/CodeGenOpenACC/atomic-capture.cpp b/clang/test/CIR/CodeGenOpenACC/atomic-capture.cpp
index 8bdffb41d1890..145c04268805f 100644
--- a/clang/test/CIR/CodeGenOpenACC/atomic-capture.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/atomic-capture.cpp
@@ -23,6 +23,7 @@ void use(int x, int v, float f, HasOps ops) {
// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(ne, %[[X_LOAD]], %[[V_LOAD]]) : !s32i, !cir.bool
// CHECK-NEXT: %[[IF_COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP:.*]] : !cir.bool to i1
// CHECK-NEXT: acc.atomic.capture if(%[[IF_COND_CAST]]) {
+ // CHECK-NEXT: acc.atomic.read %[[V_ALLOCA]] = %[[X_ALLOCA]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !s32i
// CHECK-NEXT: acc.atomic.update %[[X_ALLOCA]] : !cir.ptr<!s32i> {
// CHECK-NEXT: ^bb0(%[[X_VAR:.*]]: !s32i{{.*}}):
// CHECK-NEXT: %[[X_VAR_ALLOC:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x_var", init]
@@ -35,7 +36,6 @@ void use(int x, int v, float f, HasOps ops) {
// CHECK-NEXT: %[[X_VAR_LOAD:.*]] = cir.load{{.*}} %[[X_VAR_ALLOC]] : !cir.ptr<!s32i>, !s32i
// CHECK-NEXT: acc.yield %[[X_VAR_LOAD]] : !s32i
// CHECK-NEXT: }
- // CHECK-NEXT: acc.atomic.read %[[V_ALLOCA]] = %[[X_ALLOCA]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !s32i
// CHECK-NEXT: }
#pragma acc atomic capture if (x != v)
v = x++;
@@ -59,6 +59,7 @@ void use(int x, int v, float f, HasOps ops) {
v = ++x;
// CHECK-NEXT: acc.atomic.capture {
+ // CHECK-NEXT: acc.atomic.read %[[V_ALLOCA]] = %[[X_ALLOCA]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !s32i
// CHECK-NEXT: acc.atomic.update %[[X_ALLOCA]] : !cir.ptr<!s32i> {
// CHECK-NEXT: ^bb0(%[[X_VAR:.*]]: !s32i{{.*}}):
// CHECK-NEXT: %[[X_VAR_ALLOC:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x_var", init]
@@ -71,7 +72,6 @@ void use(int x, int v, float f, HasOps ops) {
// CHECK-NEXT: %[[X_VAR_LOAD:.*]] = cir.load{{.*}} %[[X_VAR_ALLOC]] : !cir.ptr<!s32i>, !s32i
// CHECK-NEXT: acc.yield %[[X_VAR_LOAD]] : !s32i
// CHECK-NEXT: }
- // CHECK-NEXT: acc.atomic.read %[[V_ALLOCA]] = %[[X_ALLOCA]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>, !s32i
// CHECK-NEXT: }
#pragma acc atomic capture
v = x--;
``````````
</details>
https://github.com/llvm/llvm-project/pull/168717
More information about the cfe-commits
mailing list