[clang] b307347 - [OpenACC][CIR] Lowering for atomic-read (#164299)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Oct 22 06:10:50 PDT 2025
Author: Erich Keane
Date: 2025-10-22T06:10:46-07:00
New Revision: b3073470424c9ef4c5f319d3eed4d42170e15cf1
URL: https://github.com/llvm/llvm-project/commit/b3073470424c9ef4c5f319d3eed4d42170e15cf1
DIFF: https://github.com/llvm/llvm-project/commit/b3073470424c9ef4c5f319d3eed4d42170e15cf1.diff
LOG: [OpenACC][CIR] Lowering for atomic-read (#164299)
The OpenACC spec allows only `v = x` form for atomic-read, and only when
both are L-values. The result is this ends up being a pretty trivial
patch, however it adds a decent amount of infrastructure for the other
forms of atomic.
Additionally, the 3.4 spec starts allowing the 'if' clause on atomic,
which has recently been added to the ACC dialect. This patch also
ensures that can be lowered as well. Extensive testing of this feature
was done on other clauses, so there isn't much further work/testing to
be done for it.
Added:
clang/test/CIR/CodeGenOpenACC/atomic-read.cpp
Modified:
clang/include/clang/AST/StmtOpenACC.h
clang/lib/AST/StmtOpenACC.cpp
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h
index 8b4554e996326..4d52805033410 100644
--- a/clang/include/clang/AST/StmtOpenACC.h
+++ b/clang/include/clang/AST/StmtOpenACC.h
@@ -815,6 +815,17 @@ class OpenACCAtomicConstruct final
Stmt *getAssociatedStmt() {
return OpenACCAssociatedStmtConstruct::getAssociatedStmt();
}
+
+ // A struct to represent a broken-down version of the associated statement,
+ // providing the information specified in OpenACC3.3 Section 2.12.
+ struct StmtInfo {
+ const Expr *V;
+ const Expr *X;
+ // TODO: OpenACC: We should expand this as we're implementing the other
+ // atomic construct kinds.
+ };
+
+ const StmtInfo getAssociatedStmtInfo() const;
};
} // namespace clang
diff --git a/clang/lib/AST/StmtOpenACC.cpp b/clang/lib/AST/StmtOpenACC.cpp
index 07e3de8eeb00d..2b56c1eea547c 100644
--- a/clang/lib/AST/StmtOpenACC.cpp
+++ b/clang/lib/AST/StmtOpenACC.cpp
@@ -12,7 +12,9 @@
#include "clang/AST/StmtOpenACC.h"
#include "clang/AST/ASTContext.h"
+#include "clang/AST/ExprCXX.h"
#include "clang/AST/StmtCXX.h"
+
using namespace clang;
OpenACCComputeConstruct *
@@ -322,6 +324,38 @@ OpenACCAtomicConstruct *OpenACCAtomicConstruct::Create(
return Inst;
}
+const OpenACCAtomicConstruct::StmtInfo
+OpenACCAtomicConstruct::getAssociatedStmtInfo() const {
+ // This ends up being a vastly simplified version of SemaOpenACCAtomic, since
+ // it doesn't have to worry about erroring out, but we should do a lot of
+ // asserts to ensure we don't get off into the weeds.
+ assert(getAssociatedStmt() && "invalid associated stmt?");
+
+ switch (AtomicKind) {
+ case OpenACCAtomicKind::None:
+ case OpenACCAtomicKind::Write:
+ case OpenACCAtomicKind::Update:
+ case OpenACCAtomicKind::Capture:
+ assert(false && "Only 'read' has been implemented here");
+ return {};
+ case OpenACCAtomicKind::Read: {
+ // Read only supports the format 'v = x'; where both sides are a scalar
+ // expression. This can come in 2 forms; BinaryOperator or
+ // CXXOperatorCallExpr (rarely).
+ const Expr *AssignExpr = cast<const Expr>(getAssociatedStmt());
+ if (const auto *BO = dyn_cast<BinaryOperator>(AssignExpr)) {
+ assert(BO->getOpcode() == BO_Assign);
+ return {BO->getLHS()->IgnoreImpCasts(), BO->getRHS()->IgnoreImpCasts()};
+ }
+
+ const auto *OO = cast<CXXOperatorCallExpr>(AssignExpr);
+ assert(OO->getOperator() == OO_Equal);
+
+ return {OO->getArg(0)->IgnoreImpCasts(), OO->getArg(1)->IgnoreImpCasts()};
+ }
+ }
+}
+
OpenACCCacheConstruct *OpenACCCacheConstruct::CreateEmpty(const ASTContext &C,
unsigned NumVars) {
void *Mem =
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index ce4ae7ec5efc4..385f89c5544d6 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -553,12 +553,15 @@ class OpenACCClauseCIREmitter final
}
void VisitIfClause(const OpenACCIfClause &clause) {
- if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
- mlir::acc::KernelsOp, mlir::acc::InitOp,
- mlir::acc::ShutdownOp, mlir::acc::SetOp,
- mlir::acc::DataOp, mlir::acc::WaitOp,
- mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
- mlir::acc::ExitDataOp, mlir::acc::UpdateOp>) {
+ if constexpr (isOneOfTypes<
+ OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
+ mlir::acc::KernelsOp, mlir::acc::InitOp,
+ mlir::acc::ShutdownOp, mlir::acc::SetOp,
+ mlir::acc::DataOp, mlir::acc::WaitOp,
+ mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
+ mlir::acc::ExitDataOp, mlir::acc::UpdateOp,
+ mlir::acc::AtomicReadOp, mlir::acc::AtomicWriteOp,
+ mlir::acc::AtomicUpdateOp, mlir::acc::AtomicCaptureOp>) {
operation.getIfCondMutable().append(
createCondition(clause.getConditionExpr()));
} else if constexpr (isCombinedType<OpTy>) {
@@ -1144,6 +1147,10 @@ EXPL_SPEC(mlir::acc::HostDataOp)
EXPL_SPEC(mlir::acc::EnterDataOp)
EXPL_SPEC(mlir::acc::ExitDataOp)
EXPL_SPEC(mlir::acc::UpdateOp)
+EXPL_SPEC(mlir::acc::AtomicReadOp)
+EXPL_SPEC(mlir::acc::AtomicWriteOp)
+EXPL_SPEC(mlir::acc::AtomicCaptureOp)
+EXPL_SPEC(mlir::acc::AtomicUpdateOp)
#undef EXPL_SPEC
template <typename ComputeOp, typename LoopOp>
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index e89393c92db33..02bb46d0e4466 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -306,6 +306,29 @@ CIRGenFunction::emitOpenACCCacheConstruct(const OpenACCCacheConstruct &s) {
mlir::LogicalResult
CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {
- cgm.errorNYI(s.getSourceRange(), "OpenACC Atomic Construct");
- return mlir::failure();
+ // For now, we are only support 'read', so diagnose. We can switch on the kind
+ // later once we start implementing the other 3 forms.
+ if (s.getAtomicKind() != OpenACCAtomicKind::Read) {
+ cgm.errorNYI(s.getSourceRange(), "OpenACC Atomic Construct");
+ return mlir::failure();
+ }
+
+ // While Atomic is an 'associated statement' construct, it 'steals' the
+ // expression it is associated with rather than emitting it inside of it. So
+ // it has custom emit logic.
+ mlir::Location start = getLoc(s.getSourceRange().getBegin());
+ OpenACCAtomicConstruct::StmtInfo inf = s.getAssociatedStmtInfo();
+ // Atomic 'read' only permits 'v = x', where v and x are both scalar L values.
+ // The getAssociatedStmtInfo strips off implicit casts, which includes
+ // implicit conversions and L-to-R-Value conversions, so we can just emit it
+ // as an L value. The Flang implementation has no problem with
diff erent
+ // types, so it appears that the dialect can handle the conversions.
+ mlir::Value v = emitLValue(inf.V).getPointer();
+ mlir::Value x = emitLValue(inf.X).getPointer();
+ mlir::Type resTy = convertType(inf.V->getType());
+ auto op = mlir::acc::AtomicReadOp::create(builder, start, x, v, resTy,
+ /*ifCond=*/{});
+ emitOpenACCClauses(op, s.getDirectiveKind(), s.getDirectiveLoc(),
+ s.clauses());
+ return mlir::success();
}
diff --git a/clang/test/CIR/CodeGenOpenACC/atomic-read.cpp b/clang/test/CIR/CodeGenOpenACC/atomic-read.cpp
new file mode 100644
index 0000000000000..9882f050045d3
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/atomic-read.cpp
@@ -0,0 +1,24 @@
+// 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
+
+void use(int x, unsigned int y, float f) {
+ // CHECK: cir.func{{.*}}(%[[X_ARG:.*]]: !s32i{{.*}}, %[[Y_ARG:.*]]: !u32i{{.*}}, %[[F_ARG:.*]]: !cir.float{{.*}}){{.*}}{
+ // CHECK-NEXT: %[[X_ALLOC:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["x", init]
+ // CHECK-NEXT: %[[Y_ALLOC:.*]] = cir.alloca !u32i, !cir.ptr<!u32i>, ["y", init]
+ // CHECK-NEXT: %[[F_ALLOC:.*]] = cir.alloca !cir.float, !cir.ptr<!cir.float>, ["f", init]
+ // CHECK-NEXT: cir.store %[[X_ARG]], %[[X_ALLOC]] : !s32i, !cir.ptr<!s32i>
+ // CHECK-NEXT: cir.store %[[Y_ARG]], %[[Y_ALLOC]] : !u32i, !cir.ptr<!u32i>
+ // CHECK-NEXT: cir.store %[[F_ARG]], %[[F_ALLOC]] : !cir.float, !cir.ptr<!cir.float>
+
+ // CHECK-NEXT: acc.atomic.read %[[X_ALLOC]] = %[[Y_ALLOC]] : !cir.ptr<!s32i>, !cir.ptr<!u32i>, !s32i
+#pragma acc atomic read
+ x = y;
+
+ // CHECK-NEXT: %[[X_LOAD:.*]] = cir.load{{.*}} %[[X_ALLOC]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[X_CAST:.*]] = cir.cast integral %[[X_LOAD]] : !s32i -> !u32i
+ // CHECK-NEXT: %[[Y_LOAD:.*]] = cir.load{{.*}} %[[Y_ALLOC]] : !cir.ptr<!u32i>, !u32i
+ // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[X_CAST]], %[[Y_LOAD]]) : !u32i, !cir.bool
+ // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]] : !cir.bool to i1
+ // CHECK-NEXT: acc.atomic.read if(%[[CMP_CAST]]) %[[F_ALLOC]] = %[[Y_ALLOC]] : !cir.ptr<!cir.float>, !cir.ptr<!u32i>, !cir.float
+#pragma acc atomic read if (x == y)
+ f = y;
+}
More information about the cfe-commits
mailing list