[OpenACC][CIR] 'declare' lowering for globals/ns/struct-scopes (+create) #169409
+469
−52
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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.
erichkeane
requested review from
andykaylor,
bcardosolopes,
lanza and
xlauko
as code owners
llvmbot
added
clang
Member
|
@llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) ChangesThis 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. Patch is 36.37 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/169409.diff 8 Files Affected:
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_GLOB...
[truncated]
|
Member
|
@llvm/pr-subscribers-clangir Author: Erich Keane (erichkeane) ChangesThis 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. Patch is 36.37 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/169409.diff 8 Files Affected:
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_GLOB...
[truncated]
|
erichkeane
commented
Nov 24, 2025
andykaylor
approved these changes
Nov 24, 2025
Contributor
andykaylor
left a comment
andykaylor
left a comment
There was a problem hiding this comment.
Looks good to me, with just one nit.
| while (const auto *ase = dyn_cast<ArraySectionExpr>(curVarExpr)) | ||
| curVarExpr = ase->getBase()->IgnoreParenImpCasts(); | ||
|
|
||
| if (const auto *DRE = dyn_cast<DeclRefExpr>(curVarExpr)) |
Contributor
There was a problem hiding this comment.
Suggested change
| if (const auto *DRE = dyn_cast<DeclRefExpr>(curVarExpr)) | |
| if (const auto *dre= dyn_cast<DeclRefExpr>(curVarExpr)) |
🐧 Linux x64 Test Results
|
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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.