[clang] [OpenACC][CIR] 'declare' lowering for globals/ns/struct-scopes (+create) (PR #169409)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Mon Nov 24 20:07:51 PST 2025
https://github.com/erichkeane updated https://github.com/llvm/llvm-project/pull/169409
>From f8a35329c552384dae0318013b811a95f86e1fe3 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Mon, 24 Nov 2025 11:26:37 -0800
Subject: [PATCH 1/4] [OpenACC][CIR] 'declare' lowering for
globals/ns/struct-scopes (+create)
This patch does the lowering for a 'declare' construct that is not a
function-local-scope. It also does the lowering for 'create', which has
an entry-op of create and exit-op of delete.
Global/NS/Struct scope 'declare's emit a single 'acc_ctor' and 'acc_dtor'
(except in the case of 'link') per variable referenced. The ctor is the
entry op followed by a declare_enter. The dtor is a get_device_ptr,
followed by a declare_exit, followed by a delete(exit op). This DOES
include any necessary bounds.
This patch implements all of the above. We use a separate 'visitor' for
the clauses here since it is particularly different from the other uses,
AND there are only 4 valid clauses. Additionally, we had to split the
modifier conversion into its own 'helpers' file, which will hopefully
get some additional use in the future.
---
clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 165 +++++++++++-
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 6 +-
clang/lib/CIR/CodeGen/CIRGenModule.h | 8 +
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 34 +--
clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h | 43 +++
.../CIR/CodeGenOpenACC/declare-create.cpp | 253 +++++++++++++++++-
.../openacc-not-implemented-global.cpp | 6 +-
.../openacc-not-implemented.cpp | 5 -
8 files changed, 468 insertions(+), 52 deletions(-)
create mode 100644 clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h
delete mode 100644 clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index 41a193e4d85c5..759eef2f378f8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -11,8 +11,11 @@
//===----------------------------------------------------------------------===//
#include "CIRGenFunction.h"
+#include "CIRGenOpenACCHelpers.h"
+
#include "mlir/Dialect/OpenACC/OpenACC.h"
#include "clang/AST/DeclOpenACC.h"
+#include "llvm/Support/SaveAndRestore.h"
using namespace clang;
using namespace clang::CIRGen;
@@ -96,6 +99,13 @@ struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup {
};
} // namespace
+void CIRGenModule::emitGlobalOpenACCDecl(const OpenACCConstructDecl *d) {
+ if (const auto *rd = dyn_cast<OpenACCRoutineDecl>(d))
+ emitGlobalOpenACCRoutineDecl(rd);
+ else
+ emitGlobalOpenACCDeclareDecl(cast<OpenACCDeclareDecl>(d));
+}
+
void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) {
mlir::Location exprLoc = cgm.getLoc(d.getBeginLoc());
auto enterOp = mlir::acc::DeclareEnterOp::create(
@@ -109,15 +119,156 @@ void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) {
enterOp);
}
+// Helper function that gets the declaration referenced by the declare clause.
+// This is a simplified verison of the work that `getOpenACCDataOperandInfo`
+// does, as it only has to get forms that 'declare' does.
+static const Decl *getDeclareReferencedDecl(const Expr *e) {
+ const Expr *curVarExpr = e->IgnoreParenImpCasts();
+
+ // Since we allow array sections, we have to unpack the array sections here.
+ // We don't have to worry about other bounds, since only variable or array
+ // name (plus array sections as an extension) are permitted.
+ while (const auto *ase = dyn_cast<ArraySectionExpr>(curVarExpr))
+ curVarExpr = ase->getBase()->IgnoreParenImpCasts();
+
+ if (const auto *DRE = dyn_cast<DeclRefExpr>(curVarExpr))
+ return DRE->getFoundDecl()->getCanonicalDecl();
+
+ // MemberExpr is allowed when it is implicit 'this'.
+ return cast<MemberExpr>(curVarExpr)->getMemberDecl()->getCanonicalDecl();
+}
+
+template <typename BeforeOpTy, typename DataClauseTy>
+void CIRGenModule::emitGlobalOpenACCDeclareDataOperands(
+ const Expr *varOperand, DataClauseTy dataClause,
+ OpenACCModifierKind modifiers, bool structured, bool implicit,
+ bool requiresDtor) {
+ // This is a template argument so that we don't have to include all of
+ // mlir::acc into CIRGenModule.
+ static_assert(std::is_same_v<DataClauseTy, mlir::acc::DataClause>);
+ mlir::Location exprLoc = getLoc(varOperand->getBeginLoc());
+ const Decl *refedDecl = getDeclareReferencedDecl(varOperand);
+ StringRef varName = getMangledName(GlobalDecl{cast<VarDecl>(refedDecl)});
+
+ // We have to emit two separate functions in this case, an acc_ctor and an
+ // acc_dtor. These two sections are/should remain reasonably equal, however
+ // the order of the clauses/vs-enter&exit in them makes combining these two
+ // sections not particularly attractive, so we have a bit of repetition.
+ {
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ auto ctorOp = mlir::acc::GlobalConstructorOp::create(
+ builder, exprLoc, (varName + "_acc_ctor").str());
+ getModule().push_back(ctorOp);
+ mlir::Block *block = builder.createBlock(&ctorOp.getRegion(),
+ ctorOp.getRegion().end(), {}, {});
+ builder.setInsertionPointToEnd(block);
+ // These things are close enough to a function handling-wise we can just
+ // create this here.
+ CIRGenFunction cgf{*this, builder, true};
+ llvm::SaveAndRestore<CIRGenFunction *> savedCGF(curCGF, &cgf);
+ cgf.curFn = ctorOp;
+ CIRGenFunction::SourceLocRAIIObject fnLoc{cgf, exprLoc};
+
+ // This gets the information we need, PLUS emits the bounds correctly, so we
+ // have to do this in both enter and exit.
+ CIRGenFunction::OpenACCDataOperandInfo inf =
+ cgf.getOpenACCDataOperandInfo(varOperand);
+ auto beforeOp =
+ BeforeOpTy::create(builder, exprLoc, inf.varValue, structured, implicit,
+ inf.name, inf.bounds);
+ beforeOp.setDataClause(dataClause);
+ beforeOp.setModifiers(convertOpenACCModifiers(modifiers));
+
+ mlir::acc::DeclareEnterOp::create(
+ builder, exprLoc, mlir::acc::DeclareTokenType::get(&getMLIRContext()),
+ beforeOp.getResult());
+
+ mlir::acc::TerminatorOp::create(builder, exprLoc);
+ }
+
+ // copyin, create, and device_resident require a destructor, link does not. In
+ // the case of the first three, they are all a 'getdeviceptr', followed by the
+ // declare_exit, followed by a delete op in the destructor region.
+ if (requiresDtor) {
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ auto ctorOp = mlir::acc::GlobalDestructorOp::create(
+ builder, exprLoc, (varName + "_acc_dtor").str());
+ getModule().push_back(ctorOp);
+ mlir::Block *block = builder.createBlock(&ctorOp.getRegion(),
+ ctorOp.getRegion().end(), {}, {});
+ builder.setInsertionPointToEnd(block);
+
+ // These things are close enough to a function handling-wise we can just
+ // create this here.
+ CIRGenFunction cgf{*this, builder, true};
+ llvm::SaveAndRestore<CIRGenFunction *> savedCGF(curCGF, &cgf);
+ cgf.curFn = ctorOp;
+ CIRGenFunction::SourceLocRAIIObject fnLoc{cgf, exprLoc};
+
+ CIRGenFunction::OpenACCDataOperandInfo inf =
+ cgf.getOpenACCDataOperandInfo(varOperand);
+ auto getDevPtr = mlir::acc::GetDevicePtrOp::create(
+ builder, exprLoc, inf.varValue, structured, implicit, inf.name,
+ inf.bounds);
+ getDevPtr.setDataClause(dataClause);
+ getDevPtr.setModifiers(convertOpenACCModifiers(modifiers));
+
+ mlir::acc::DeclareExitOp::create(builder, exprLoc, /*token=*/mlir::Value{},
+ getDevPtr.getResult());
+ auto deleteOp = mlir::acc::DeleteOp::create(
+ builder, exprLoc, getDevPtr, structured, implicit, inf.name, {});
+ deleteOp.setDataClause(dataClause);
+ deleteOp.setModifiers(convertOpenACCModifiers(modifiers));
+ mlir::acc::TerminatorOp::create(builder, exprLoc);
+ }
+}
+namespace {
+// This class emits all of the information for a 'declare' at a global/ns/class
+// scope. Each clause results in its own acc_ctor and acc_dtor for the variable.
+// This class creates those and emits them properly.
+// This behavior is unique/special enough from the emission of statement-level
+// clauses that it doesn't really make sense to use that clause visitor.
+class OpenACCGlobalDeclareClauseEmitter final
+ : public OpenACCClauseVisitor<OpenACCGlobalDeclareClauseEmitter> {
+ CIRGenModule &cgm;
+ void clauseNotImplemented(const OpenACCClause &c) {
+ cgm.errorNYI(c.getSourceRange(), "OpenACC Global Declare Clause",
+ c.getClauseKind());
+ }
+
+public:
+ OpenACCGlobalDeclareClauseEmitter(CIRGenModule &cgm) : cgm(cgm) {}
+
+ void VisitClause(const OpenACCClause &clause) {
+ clauseNotImplemented(clause);
+ }
+
+ void emitClauses(ArrayRef<const OpenACCClause *> clauses) {
+ this->VisitClauseList(clauses);
+ }
+
+ void VisitCreateClause(const OpenACCCreateClause &clause) {
+ for (const Expr *var : clause.getVarList())
+ cgm.emitGlobalOpenACCDeclareDataOperands<mlir::acc::CreateOp>(
+ var, mlir::acc::DataClause::acc_create, {}, /*structured=*/true,
+ /*implicit=*/false, /*requiresDtor=*/true);
+ }
+};
+} // namespace
+
+void CIRGenModule::emitGlobalOpenACCDeclareDecl(const OpenACCDeclareDecl *d) {
+ // Declare creates 1 'acc_ctor' and 0-1 'acc_dtor' per clause, since it needs
+ // a unique one on a per-variable basis. We can just use a clause emitter to
+ // do all the work.
+ mlir::OpBuilder::InsertionGuard guardCase(builder);
+ OpenACCGlobalDeclareClauseEmitter em{*this};
+ em.emitClauses(d->clauses());
+}
+
void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) {
getCIRGenModule().errorNYI(d.getSourceRange(), "OpenACC Routine Construct");
}
-void CIRGenModule::emitGlobalOpenACCDecl(const OpenACCConstructDecl *d) {
- if (isa<OpenACCRoutineDecl>(d))
- errorNYI(d->getSourceRange(), "OpenACC Routine Construct");
- else if (isa<OpenACCDeclareDecl>(d))
- errorNYI(d->getSourceRange(), "OpenACC Declare Construct");
- else
- llvm_unreachable("unknown OpenACC declaration kind?");
+void CIRGenModule::emitGlobalOpenACCRoutineDecl(const OpenACCRoutineDecl *d) {
+ errorNYI(d->getSourceRange(), "OpenACC Global Routine Construct");
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 251c99c8cd45b..809c24f8aa670 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -1513,10 +1513,10 @@ void CIRGenModule::emitTopLevelDecl(Decl *decl) {
break;
}
case Decl::OpenACCRoutine:
- emitGlobalOpenACCDecl(cast<OpenACCRoutineDecl>(decl));
+ emitGlobalOpenACCRoutineDecl(cast<OpenACCRoutineDecl>(decl));
break;
case Decl::OpenACCDeclare:
- emitGlobalOpenACCDecl(cast<OpenACCDeclareDecl>(decl));
+ emitGlobalOpenACCDeclareDecl(cast<OpenACCDeclareDecl>(decl));
break;
case Decl::Enum:
case Decl::Using: // using X; [C++]
@@ -1560,7 +1560,7 @@ void CIRGenModule::emitTopLevelDecl(Decl *decl) {
CXXRecordDecl *crd = cast<CXXRecordDecl>(decl);
assert(!cir::MissingFeatures::generateDebugInfo());
for (auto *childDecl : crd->decls())
- if (isa<VarDecl, CXXRecordDecl, EnumDecl>(childDecl))
+ if (isa<VarDecl, CXXRecordDecl, EnumDecl, OpenACCDeclareDecl>(childDecl))
emitTopLevelDecl(childDecl);
break;
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h
index 2c45bb238f95a..6600d086f8f61 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -453,6 +453,14 @@ class CIRGenModule : public CIRGenTypeCache {
bool performInit);
void emitGlobalOpenACCDecl(const clang::OpenACCConstructDecl *cd);
+ void emitGlobalOpenACCRoutineDecl(const clang::OpenACCRoutineDecl *cd);
+ void emitGlobalOpenACCDeclareDecl(const clang::OpenACCDeclareDecl *cd);
+ template <typename BeforeOpTy, typename DataClauseTy>
+ void emitGlobalOpenACCDeclareDataOperands(const Expr *varOperand,
+ DataClauseTy dataClause,
+ OpenACCModifierKind modifiers,
+ bool structured, bool implicit,
+ bool requiresDtor);
// C++ related functions.
void emitDeclContext(const DeclContext *dc);
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 60a089fe0e936..25ba6b0369bce 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -14,6 +14,7 @@
#include "CIRGenCXXABI.h"
#include "CIRGenFunction.h"
+#include "CIRGenOpenACCHelpers.h"
#include "CIRGenOpenACCRecipe.h"
#include "clang/AST/ExprCXX.h"
@@ -182,33 +183,6 @@ class OpenACCClauseCIREmitter final
dataOperands.append(computeEmitter.dataOperands);
}
- mlir::acc::DataClauseModifier
- convertModifiers(OpenACCModifierKind modifiers) {
- using namespace mlir::acc;
- static_assert(static_cast<int>(OpenACCModifierKind::Zero) ==
- static_cast<int>(DataClauseModifier::zero) &&
- static_cast<int>(OpenACCModifierKind::Readonly) ==
- static_cast<int>(DataClauseModifier::readonly) &&
- static_cast<int>(OpenACCModifierKind::AlwaysIn) ==
- static_cast<int>(DataClauseModifier::alwaysin) &&
- static_cast<int>(OpenACCModifierKind::AlwaysOut) ==
- static_cast<int>(DataClauseModifier::alwaysout) &&
- static_cast<int>(OpenACCModifierKind::Capture) ==
- static_cast<int>(DataClauseModifier::capture));
-
- DataClauseModifier mlirModifiers{};
-
- // The MLIR representation of this represents `always` as `alwaysin` +
- // `alwaysout`. So do a small fixup here.
- if (isOpenACCModifierBitSet(modifiers, OpenACCModifierKind::Always)) {
- mlirModifiers = mlirModifiers | DataClauseModifier::always;
- modifiers &= ~OpenACCModifierKind::Always;
- }
-
- mlirModifiers = mlirModifiers | static_cast<DataClauseModifier>(modifiers);
- return mlirModifiers;
- }
-
template <typename BeforeOpTy, typename AfterOpTy>
void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
OpenACCModifierKind modifiers, bool structured,
@@ -243,8 +217,8 @@ class OpenACCClauseCIREmitter final
// Set the 'rest' of the info for both operations.
beforeOp.setDataClause(dataClause);
afterOp.setDataClause(dataClause);
- beforeOp.setModifiers(convertModifiers(modifiers));
- afterOp.setModifiers(convertModifiers(modifiers));
+ beforeOp.setModifiers(convertOpenACCModifiers(modifiers));
+ afterOp.setModifiers(convertOpenACCModifiers(modifiers));
// Make sure we record these, so 'async' values can be updated later.
dataOperands.push_back(beforeOp.getOperation());
@@ -264,7 +238,7 @@ class OpenACCClauseCIREmitter final
// Set the 'rest' of the info for the operation.
beforeOp.setDataClause(dataClause);
- beforeOp.setModifiers(convertModifiers(modifiers));
+ beforeOp.setModifiers(convertOpenACCModifiers(modifiers));
// Make sure we record these, so 'async' values can be updated later.
dataOperands.push_back(beforeOp.getOperation());
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h b/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h
new file mode 100644
index 0000000000000..5bcc9f57d67b1
--- /dev/null
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCHelpers.h
@@ -0,0 +1,43 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This contains helpers for OpenACC emission that don't need to be in
+// CIRGenModule, but can't live in a single .cpp file.
+//
+//===----------------------------------------------------------------------===//
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "clang/AST/DeclOpenACC.h"
+
+namespace clang::CIRGen {
+inline mlir::acc::DataClauseModifier
+convertOpenACCModifiers(OpenACCModifierKind modifiers) {
+ using namespace mlir::acc;
+ static_assert(static_cast<int>(OpenACCModifierKind::Zero) ==
+ static_cast<int>(DataClauseModifier::zero) &&
+ static_cast<int>(OpenACCModifierKind::Readonly) ==
+ static_cast<int>(DataClauseModifier::readonly) &&
+ static_cast<int>(OpenACCModifierKind::AlwaysIn) ==
+ static_cast<int>(DataClauseModifier::alwaysin) &&
+ static_cast<int>(OpenACCModifierKind::AlwaysOut) ==
+ static_cast<int>(DataClauseModifier::alwaysout) &&
+ static_cast<int>(OpenACCModifierKind::Capture) ==
+ static_cast<int>(DataClauseModifier::capture));
+
+ DataClauseModifier mlirModifiers{};
+
+ // The MLIR representation of this represents `always` as `alwaysin` +
+ // `alwaysout`. So do a small fixup here.
+ if (isOpenACCModifierBitSet(modifiers, OpenACCModifierKind::Always)) {
+ mlirModifiers = mlirModifiers | DataClauseModifier::always;
+ modifiers &= ~OpenACCModifierKind::Always;
+ }
+
+ mlirModifiers = mlirModifiers | static_cast<DataClauseModifier>(modifiers);
+ return mlirModifiers;
+}
+} // namespace clang::CIRGen
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-create.cpp b/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
index ef2f1de19ea96..988454ea8a3d2 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
@@ -5,14 +5,259 @@ struct HasSideEffects {
~HasSideEffects();
};
-// TODO: OpenACC: Implement 'global', NS lowering.
+HasSideEffects GlobalHSE1;
+HasSideEffects GlobalHSEArr[5];
+int GlobalInt1;
+
+#pragma acc declare create(GlobalHSE1, GlobalInt1, GlobalHSEArr[1:1])
+// CHECK: acc.global_ctor @GlobalHSE1_acc_ctor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSE1 : !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "GlobalHSE1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @GlobalHSE1_acc_dtor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSE1 : !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_create>, name = "GlobalHSE1"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, name = "GlobalHSE1"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+//
+// CHECK: acc.global_ctor @GlobalInt1_acc_ctor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalInt1 : !cir.ptr<!s32i>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "GlobalInt1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @GlobalInt1_acc_dtor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalInt1 : !cir.ptr<!s32i>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_create>, name = "GlobalInt1"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "GlobalInt1"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+//
+// CHECK: acc.global_ctor @GlobalHSEArr_acc_ctor {
+// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
+// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
+// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
+// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
+// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "GlobalHSEArr[1:1]"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @GlobalHSEArr_acc_dtor {
+// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
+// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
+// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
+// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
+// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_create>, name = "GlobalHSEArr[1:1]"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_create>, name = "GlobalHSEArr[1:1]"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+
+namespace NS {
+
+HasSideEffects NSHSE1;
+HasSideEffects NSHSEArr[5];
+int NSInt1;
+
+#pragma acc declare create(zero: NSHSE1, NSInt1, NSHSEArr[1:1])
+// CHECK: acc.global_ctor @{{.*}}NSHSE1{{.*}}_acc_ctor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "NSHSE1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @{{.*}}NSHSE1{{.*}}_acc_dtor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_create>, name = "NSHSE1"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, name = "NSHSE1"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+//
+// CHECK: acc.global_ctor @{{.*}}NSInt1{{.*}}_acc_ctor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "NSInt1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @{{.*}}NSInt1{{.*}}_acc_dtor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_create>, name = "NSInt1"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "NSInt1"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+//
+// CHECK: acc.global_ctor @{{.*}}NSHSEArr{{.*}}_acc_ctor {
+// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
+// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
+// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
+// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
+// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "NSHSEArr[1:1]"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @{{.*}}NSHSEArr{{.*}}_acc_dtor {
+// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
+// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
+// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
+// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
+// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_create>, name = "NSHSEArr[1:1]"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_create>, name = "NSHSEArr[1:1]"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+
+
+} // namespace NS
+
+namespace {
+
+HasSideEffects AnonNSHSE1;
+HasSideEffects AnonNSHSEArr[5];
+int AnonNSInt1;
+
+#pragma acc declare create(AnonNSHSE1, AnonNSInt1, AnonNSHSEArr[1:1])
+// CHECK: acc.global_ctor @{{.*}}AnonNSHSE1{{.*}}_acc_ctor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "AnonNSHSE1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @{{.*}}AnonNSHSE1{{.*}}_acc_dtor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_create>, name = "AnonNSHSE1"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, name = "AnonNSHSE1"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+//
+// CHECK: acc.global_ctor @{{.*}}AnonNSInt1{{.*}}_acc_ctor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSInt1{{.*}} : !cir.ptr<!s32i>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "AnonNSInt1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @{{.*}}AnonNSInt1{{.*}}_acc_dtor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSInt1{{.*}} : !cir.ptr<!s32i>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_create>, name = "AnonNSInt1"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "AnonNSInt1"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+//
+// CHECK: acc.global_ctor @{{.*}}AnonNSHSEArr{{.*}}_acc_ctor {
+// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
+// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
+// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
+// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
+// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "AnonNSHSEArr[1:1]"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @{{.*}}AnonNSHSEArr{{.*}}_acc_dtor {
+// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
+// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
+// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
+// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
+// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_create>, name = "AnonNSHSEArr[1:1]"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_create>, name = "AnonNSHSEArr[1:1]"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+
+} // namespace NS
struct Struct {
- static const HasSideEffects StaticMemHSE;
+ static const HasSideEffects StaticMemHSE1;
static const HasSideEffects StaticMemHSEArr[5];
- static const int StaticMemInt;
+ static const int StaticMemInt1;
- // TODO: OpenACC: Implement static-local lowering.
+#pragma acc declare create(StaticMemHSE1, StaticMemInt1, StaticMemHSEArr[1:1])
+// CHECK: acc.global_ctor @{{.*}}Struct{{.*}}StaticMemHSE1{{.*}}_acc_ctor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}{{.*}}Struct{{.*}}StaticMemHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "StaticMemHSE1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @{{.*}}Struct{{.*}}StaticMemHSE1{{.*}}_acc_dtor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}Struct{{.*}}StaticMemHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_create>, name = "StaticMemHSE1"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, name = "StaticMemHSE1"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+//
+// CHECK: acc.global_ctor @{{.*}}Struct{{.*}}StaticMemInt1{{.*}}_acc_ctor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}Struct{{.*}}StaticMemInt1{{.*}} : !cir.ptr<!s32i>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "StaticMemInt1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @{{.*}}Struct{{.*}}StaticMemInt1{{.*}}_acc_dtor {
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}Struct{{.*}}StaticMemInt1{{.*}} : !cir.ptr<!s32i>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_create>, name = "StaticMemInt1"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "StaticMemInt1"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+//
+// CHECK: acc.global_ctor @{{.*}}Struct{{.*}}StaticMemHSEArr{{.*}}_acc_ctor {
+// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
+// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
+// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
+// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
+// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}Struct{{.*}}StaticMemHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "StaticMemHSEArr[1:1]"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+// CHECK: acc.global_dtor @{{.*}}Struct{{.*}}StaticMemHSEArr{{.*}}_acc_dtor {
+// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
+// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
+// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
+// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
+// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
+// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
+// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}Struct{{.*}}StaticMemHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_create>, name = "StaticMemHSEArr[1:1]"}
+// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_create>, name = "StaticMemHSEArr[1:1]"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
void MemFunc1(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) {
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}})
diff --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented-global.cpp b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented-global.cpp
index 2aa32b0484f2c..a5e4694c6f5e6 100644
--- a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented-global.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented-global.cpp
@@ -1,6 +1,6 @@
// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify
// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-llvm %s -o %t-cir.ll -verify
-int Global;
-// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
-#pragma acc declare create(Global)
+void foo() {}
+// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenACC Global Routine Construct}}
+#pragma acc routine(foo) seq
diff --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
deleted file mode 100644
index 43d91f180acaf..0000000000000
--- a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
+++ /dev/null
@@ -1,5 +0,0 @@
-// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-linux-gnu -fopenacc -fclangir -emit-cir %s -o %t.cir -verify
-
-int E, A;
-// expected-error at +1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}}
-#pragma acc declare link(E) create(A)
>From 490902e80bd303e640e5a763fa3a4c6e52362416 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Mon, 24 Nov 2025 13:39:46 -0800
Subject: [PATCH 2/4] FIX modifier emission for 'create'
---
clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 3 ++-
.../test/CIR/CodeGenOpenACC/declare-create.cpp | 18 +++++++++---------
2 files changed, 11 insertions(+), 10 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index 759eef2f378f8..4e4b8ec7495a1 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -250,7 +250,8 @@ class OpenACCGlobalDeclareClauseEmitter final
void VisitCreateClause(const OpenACCCreateClause &clause) {
for (const Expr *var : clause.getVarList())
cgm.emitGlobalOpenACCDeclareDataOperands<mlir::acc::CreateOp>(
- var, mlir::acc::DataClause::acc_create, {}, /*structured=*/true,
+ var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
+ /*structured=*/true,
/*implicit=*/false, /*requiresDtor=*/true);
}
};
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-create.cpp b/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
index 988454ea8a3d2..e5cf70190b849 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-create.cpp
@@ -75,29 +75,29 @@ int NSInt1;
#pragma acc declare create(zero: NSHSE1, NSInt1, NSHSEArr[1:1])
// CHECK: acc.global_ctor @{{.*}}NSHSE1{{.*}}_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
-// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "NSHSE1"}
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {modifiers = #acc<data_clause_modifier zero>, name = "NSHSE1"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!rec_HasSideEffects>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
// CHECK: acc.global_dtor @{{.*}}NSHSE1{{.*}}_acc_dtor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
-// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_create>, name = "NSHSE1"}
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "NSHSE1"}
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
-// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, name = "NSHSE1"}
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "NSHSE1"}
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
// CHECK: acc.global_ctor @{{.*}}NSInt1{{.*}}_acc_ctor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
-// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "NSInt1"}
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "NSInt1"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!s32i>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
// CHECK: acc.global_dtor @{{.*}}NSInt1{{.*}}_acc_dtor {
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
-// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_create>, name = "NSInt1"}
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "NSInt1"}
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
-// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "NSInt1"}
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "NSInt1"}
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
//
@@ -110,7 +110,7 @@ int NSInt1;
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
-// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {name = "NSHSEArr[1:1]"}
+// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {modifiers = #acc<data_clause_modifier zero>, name = "NSHSEArr[1:1]"}
// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
@@ -123,9 +123,9 @@ int NSInt1;
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
-// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_create>, name = "NSHSEArr[1:1]"}
+// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!rec_HasSideEffects x 5>> {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "NSHSEArr[1:1]"}
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
-// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_create>, name = "NSHSEArr[1:1]"}
+// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "NSHSEArr[1:1]"}
// CHECK-NEXT: acc.terminator
// CHECK-NEXT: }
>From b9b8e2911a70b13cbbde04326bfbebbe04fcb646 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Mon, 24 Nov 2025 19:13:46 -0800
Subject: [PATCH 3/4] Fix Andy's nit
---
clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index 4e4b8ec7495a1..d10e17ed44d60 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -131,8 +131,8 @@ static const Decl *getDeclareReferencedDecl(const Expr *e) {
while (const auto *ase = dyn_cast<ArraySectionExpr>(curVarExpr))
curVarExpr = ase->getBase()->IgnoreParenImpCasts();
- if (const auto *DRE = dyn_cast<DeclRefExpr>(curVarExpr))
- return DRE->getFoundDecl()->getCanonicalDecl();
+ if (const auto *dre = dyn_cast<DeclRefExpr>(curVarExpr))
+ return dredre->getFoundDecl()->getCanonicalDecl();
// MemberExpr is allowed when it is implicit 'this'.
return cast<MemberExpr>(curVarExpr)->getMemberDecl()->getCanonicalDecl();
>From d4a0557744b5d9e2c8cd62b072e3667feaaf3457 Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Mon, 24 Nov 2025 20:07:39 -0800
Subject: [PATCH 4/4] fixed dumb typo
---
clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index d10e17ed44d60..aeb43f2c7bbed 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -132,7 +132,7 @@ static const Decl *getDeclareReferencedDecl(const Expr *e) {
curVarExpr = ase->getBase()->IgnoreParenImpCasts();
if (const auto *dre = dyn_cast<DeclRefExpr>(curVarExpr))
- return dredre->getFoundDecl()->getCanonicalDecl();
+ return dre->getFoundDecl()->getCanonicalDecl();
// MemberExpr is allowed when it is implicit 'this'.
return cast<MemberExpr>(curVarExpr)->getMemberDecl()->getCanonicalDecl();
More information about the cfe-commits
mailing list