Skip to content

[OpenACC][CIR] Handle 'declare' construct local lowering (&link clause) #168793

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
erichkeane merged 2 commits into from
Nov 21, 2025
Merged

[OpenACC][CIR] Handle 'declare' construct local lowering (&link clause) #168793

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
22caecc
[OpenACC][CIR] Handle 'declare' construct local lowering (&link clause)
erichkeane Nov 19, 2025
126313a
clang-format
erichkeane Nov 19, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
32 changes: 31 additions & 1 deletion clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down
13 changes: 13 additions & 0 deletions clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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())
Expand Down Expand Up @@ -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>
Expand Down
130 changes: 130 additions & 0 deletions clang/test/CIR/CodeGenOpenACC/declare-link.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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;
Copy link
Collaborator Author

@erichkeane erichkeane Nov 19, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'll note all the get_global stuff is because link requires extern variables in the same scope. It perhaps wasn't the best idea to start with link, but the rest are more complicated source changes, so this was a bit of a necessary evil. I promise the next ones will be better :)

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
}

5 changes: 3 additions & 2 deletions clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
#pragma acc declare create(A)
// expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Clause: create}}
#pragma acc declare link(E) create(A)
}
Loading