[clang] 5dbe83c - [OpenACC][CIR] Handle 'declare' construct local lowering (&link clause) (#168793)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Nov 21 10:48:13 PST 2025
Author: Erich Keane
Date: 2025-11-21T10:47:52-08:00
New Revision: 5dbe83c3023a795595b52e75cdfc7835882e5db1
URL: https://github.com/llvm/llvm-project/commit/5dbe83c3023a795595b52e75cdfc7835882e5db1
DIFF: https://github.com/llvm/llvm-project/commit/5dbe83c3023a795595b52e75cdfc7835882e5db1.diff
LOG: [OpenACC][CIR] Handle 'declare' construct local lowering (&link clause) (#168793)
'declare' is a declaration directive, so it can appear at 3 places:
Global/NS scope, class scope, or local scope. This patch implements ONLY
the 'local' scope lowering for 'declare'.
A 'declare' is lowered as a 'declare_enter' and 'declare_exit'
operation, plus data operands like all others. Sema restricts the form
of some of these, but they are otherwise identical.
'declare' DOES require at least 1 clause for the examples to make sense,
so this ALSO implements 'link', which is the 'simpliest' one. It is ONLY
attached to the 'declare_enter', and doesn't require any additional work
besides a very small addition to how we handle clauses.
Added:
clang/test/CIR/CodeGenOpenACC/declare-link.cpp
Modified:
clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
Removed:
################################################################################
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index b588a50aa0404..551027bb1c8eb 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -11,13 +11,41 @@
//===----------------------------------------------------------------------===//
#include "CIRGenFunction.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "clang/AST/DeclOpenACC.h"
using namespace clang;
using namespace clang::CIRGen;
+namespace {
+struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup {
+ mlir::acc::DeclareEnterOp enterOp;
+
+ OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp) : enterOp(enterOp) {}
+
+ void emit(CIRGenFunction &cgf) override {
+ mlir::acc::DeclareExitOp::create(cgf.getBuilder(), enterOp.getLoc(),
+ enterOp, {});
+
+ // TODO(OpenACC): Some clauses require that we add info about them to the
+ // DeclareExitOp. However, we don't have any of those implemented yet, so
+ // we should add infrastructure here to do that once we have one
+ // implemented.
+ }
+};
+} // namespace
+
void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) {
- getCIRGenModule().errorNYI(d.getSourceRange(), "OpenACC Declare Construct");
+ mlir::Location exprLoc = cgm.getLoc(d.getBeginLoc());
+ auto enterOp = mlir::acc::DeclareEnterOp::create(
+ builder, exprLoc, mlir::acc::DeclareTokenType::get(&cgm.getMLIRContext()),
+ {});
+
+ emitOpenACCClauses(enterOp, OpenACCDirectiveKind::Declare, d.getBeginLoc(),
+ d.clauses());
+
+ ehStack.pushCleanup<OpenACCDeclareCleanup>(CleanupKind::NormalCleanup,
+ enterOp);
}
void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) {
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 527dfd21db8a5..c7e6a256c3868 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -876,6 +876,18 @@ class OpenACCClauseCIREmitter final
}
}
+ void VisitLinkClause(const OpenACCLinkClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::DeclareEnterOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::DeclareLinkOp>(
+ var, mlir::acc::DataClause::acc_declare_link, {},
+ /*structured=*/true,
+ /*implicit=*/false);
+ } else {
+ llvm_unreachable("Unknown construct kind in VisitLinkClause");
+ }
+ }
+
void VisitDeleteClause(const OpenACCDeleteClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
for (const Expr *var : clause.getVarList())
@@ -1151,6 +1163,7 @@ EXPL_SPEC(mlir::acc::AtomicReadOp)
EXPL_SPEC(mlir::acc::AtomicWriteOp)
EXPL_SPEC(mlir::acc::AtomicCaptureOp)
EXPL_SPEC(mlir::acc::AtomicUpdateOp)
+EXPL_SPEC(mlir::acc::DeclareEnterOp)
#undef EXPL_SPEC
template <typename ComputeOp, typename LoopOp>
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-link.cpp b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp
new file mode 100644
index 0000000000000..8494a2354c7db
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/declare-link.cpp
@@ -0,0 +1,130 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+struct HasSideEffects {
+ HasSideEffects();
+ ~HasSideEffects();
+};
+
+// TODO: OpenACC: Implement 'global', NS lowering.
+
+struct Struct {
+ static const HasSideEffects StaticMemHSE;
+ static const HasSideEffects StaticMemHSEArr[5];
+ static const int StaticMemInt;
+
+ // TODO: OpenACC: Implement static-local lowering.
+
+ void MemFunc1() {
+ // CHECK: cir.func {{.*}}MemFunc1{{.*}}({{.*}}) {
+ // CHECK-NEXT: cir.alloca{{.*}}["this"
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.load
+ extern HasSideEffects LocalHSE;
+ extern HasSideEffects LocalHSEArr[5];
+ extern int LocalInt;
+#pragma acc declare link(LocalHSE, LocalInt, LocalHSEArr[1:1])
+
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt : !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt"}
+ //
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
+ // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr[1:1]"}
+ //
+ // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+ //
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+ }
+
+ void MemFunc2();
+};
+void use() {
+ Struct s;
+ s.MemFunc1();
+}
+
+void Struct::MemFunc2() {
+ // CHECK: cir.func {{.*}}MemFunc2{{.*}}({{.*}}) {
+ // CHECK-NEXT: cir.alloca{{.*}}["this"
+ // CHECK-NEXT: cir.store
+ // CHECK-NEXT: cir.load
+ extern HasSideEffects LocalHSE2;
+ extern HasSideEffects LocalHSEArr2[5];
+ extern int LocalInt2;
+
+#pragma acc declare link(LocalHSE2, LocalInt2, LocalHSEArr2[1:1])
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE2 : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE2"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt2 : !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt2"}
+ //
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
+ // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr2 : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr2[1:1]"}
+ //
+ // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+ //
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+}
+
+extern "C" void do_thing();
+
+void NormalFunc() {
+ // CHECK: cir.func {{.*}}NormalFunc{{.*}}()
+ extern HasSideEffects LocalHSE3;
+ extern HasSideEffects LocalHSEArr3[5];
+ extern int LocalInt3;
+ // CHECK-NEXT: cir.scope
+ {
+ extern HasSideEffects InnerHSE;
+#pragma acc declare link(LocalHSE3, LocalInt3, LocalHSEArr3[1:1], InnerHSE)
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @LocalHSE3 : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[HSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "LocalHSE3"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_INT:.*]] = cir.get_global @LocalInt3 : !cir.ptr<!s32i>
+ // CHECK-NEXT: %[[INT_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_INT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "LocalInt3"}
+ //
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32
+ // CHECK-NEXT: %[[ZERO:.*]] = arith.constant 0 : i64
+ // CHECK-NEXT: %[[ONE:.*]] = arith.constant 1 : i64
+ // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[ONE]] : i64) startIdx(%[[ZERO]] : i64)
+ // CHECK-NEXT: %[[GET_LOCAL_ARR:.*]] = cir.get_global @LocalHSEArr3 : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+ // CHECK-NEXT: %[[ARR_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_ARR]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "LocalHSEArr3[1:1]"}
+ //
+ // CHECK-NEXT: %[[GET_LOCAL_HSE:.*]] = cir.get_global @InnerHSE : !cir.ptr<!rec_HasSideEffects>
+ // CHECK-NEXT: %[[INNERHSE_LINK:.*]] = acc.declare_link varPtr(%[[GET_LOCAL_HSE]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "InnerHSE"}
+ //
+ // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[HSE_LINK]], %[[INT_LINK]], %[[ARR_LINK]], %[[INNERHSE_LINK]] : !cir.ptr<!rec_HasSideEffects>, !cir.ptr<!s32i>, !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>, !cir.ptr<!rec_HasSideEffects>)
+ //
+ // CHECK
+
+ do_thing();
+ // CHECK-NEXT: cir.call @do_thing
+
+ // CHECK-NEXT: acc.declare_exit token(%[[ENTER]])
+ }
+ // CHECK-NEXT: }
+
+ do_thing();
+ // CHECK-NEXT: cir.call @do_thing
+}
+
diff --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
index e85c26718acb8..c8b85a12f84e7 100644
--- a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
@@ -1,7 +1,8 @@
// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify
void HelloWorld(int *A) {
+ extern int *E;
-// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
-#pragma acc declare create(A)
+// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenACC Clause: create}}
+#pragma acc declare link(E) create(A)
}
More information about the cfe-commits
mailing list