https://github.com/erichkeane updated https://github.com/llvm/llvm-project/pull/169409
>From f8a35329c552384dae0318013b811a95f86e1fe3 Mon Sep 17 00:00:00 2001 From: erichkeane <[email protected]> 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@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare Construct}} -#pragma acc declare create(Global) +void foo() {} +// expected-error@+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@+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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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(); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
