[clang] [OpenACC][CIR] Handle 'declare' construct local lowering (&link clause) (PR #168793)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Wed Nov 19 15:47:10 PST 2025
https://github.com/erichkeane updated https://github.com/llvm/llvm-project/pull/168793
>From 22caecc11fc0d5d6113bfd6ba24f3644316e5350 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Wed, 19 Nov 2025 13:28:18 -0800
Subject: [PATCH 1/2] [OpenACC][CIR] Handle 'declare' construct local lowering
(&link clause)
'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.
---
clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 32 ++++-
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 13 ++
.../test/CIR/CodeGenOpenACC/declare-link.cpp | 130 ++++++++++++++++++
.../openacc-not-implemented.cpp | 5 +-
4 files changed, 177 insertions(+), 3 deletions(-)
create mode 100644 clang/test/CIR/CodeGenOpenACC/declare-link.cpp
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index b588a50aa0404..f6680cbaa8c78 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -12,12 +12,42 @@
#include "CIRGenFunction.h"
#include "clang/AST/DeclOpenACC.h"
+#include "mlir/Dialect/OpenACC/OpenACC.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)
}
>From 126313abdaaf4e3756e494d4b4ca313ee7a6abaa Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Wed, 19 Nov 2025 15:46:59 -0800
Subject: [PATCH 2/2] clang-format
---
clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 32 ++++++++++-----------
1 file changed, 15 insertions(+), 17 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index f6680cbaa8c78..551027bb1c8eb 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -11,37 +11,35 @@
//===----------------------------------------------------------------------===//
#include "CIRGenFunction.h"
-#include "clang/AST/DeclOpenACC.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) {}
+struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup {
+ mlir::acc::DeclareEnterOp enterOp;
- void emit(CIRGenFunction &cgf) override {
- mlir::acc::DeclareExitOp::create(cgf.getBuilder(), enterOp.getLoc(),
- enterOp, {});
+ OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp) : enterOp(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.
- }
+ 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) {
mlir::Location exprLoc = cgm.getLoc(d.getBeginLoc());
auto enterOp = mlir::acc::DeclareEnterOp::create(
- builder, exprLoc,
- mlir::acc::DeclareTokenType::get(&cgm.getMLIRContext()), {});
+ builder, exprLoc, mlir::acc::DeclareTokenType::get(&cgm.getMLIRContext()),
+ {});
emitOpenACCClauses(enterOp, OpenACCDirectiveKind::Declare, d.getBeginLoc(),
d.clauses());
More information about the cfe-commits
mailing list