https://github.com/erichkeane created https://github.com/llvm/llvm-project/pull/168717
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. >From 5e991e1cb90ae4ea53ecf9250946495d9c5aa87a Mon Sep 17 00:00:00 2001 From: erichkeane <[email protected]> Date: Wed, 19 Nov 2025 06:16:42 -0800 Subject: [PATCH] [OpenACC][CIR] Fix atomic-capture single-line-postfix 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. --- clang/include/clang/AST/StmtOpenACC.h | 12 +++++-- clang/lib/AST/StmtOpenACC.cpp | 35 ++++++++++--------- .../CIR/CodeGenOpenACC/atomic-capture.cpp | 4 +-- 3 files changed, 29 insertions(+), 22 deletions(-) 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--; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
