diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp index 41a193e4d85c5..aeb43f2c7bbed 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(d)) + emitGlobalOpenACCRoutineDecl(rd); + else + emitGlobalOpenACCDeclareDecl(cast(d)); +} + void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) { mlir::Location exprLoc = cgm.getLoc(d.getBeginLoc()); auto enterOp = mlir::acc::DeclareEnterOp::create( @@ -109,15 +119,157 @@ 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(curVarExpr)) + curVarExpr = ase->getBase()->IgnoreParenImpCasts(); + + if (const auto *dre = dyn_cast(curVarExpr)) + return dre->getFoundDecl()->getCanonicalDecl(); + + // MemberExpr is allowed when it is implicit 'this'. + return cast(curVarExpr)->getMemberDecl()->getCanonicalDecl(); +} + +template +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); + mlir::Location exprLoc = getLoc(varOperand->getBeginLoc()); + const Decl *refedDecl = getDeclareReferencedDecl(varOperand); + StringRef varName = getMangledName(GlobalDecl{cast(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 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 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 { + 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 clauses) { + this->VisitClauseList(clauses); + } + + void VisitCreateClause(const OpenACCCreateClause &clause) { + for (const Expr *var : clause.getVarList()) + cgm.emitGlobalOpenACCDeclareDataOperands( + var, mlir::acc::DataClause::acc_create, clause.getModifierList(), + /*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(d)) - errorNYI(d->getSourceRange(), "OpenACC Routine Construct"); - else if (isa(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(decl)); + emitGlobalOpenACCRoutineDecl(cast(decl)); break; case Decl::OpenACCDeclare: - emitGlobalOpenACCDecl(cast(decl)); + emitGlobalOpenACCDeclareDecl(cast(decl)); break; case Decl::Enum: case Decl::Using: // using X; [C++] @@ -1560,7 +1560,7 @@ void CIRGenModule::emitTopLevelDecl(Decl *decl) { CXXRecordDecl *crd = cast(decl); assert(!cir::MissingFeatures::generateDebugInfo()); for (auto *childDecl : crd->decls()) - if (isa(childDecl)) + if (isa(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 + 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(OpenACCModifierKind::Zero) == - static_cast(DataClauseModifier::zero) && - static_cast(OpenACCModifierKind::Readonly) == - static_cast(DataClauseModifier::readonly) && - static_cast(OpenACCModifierKind::AlwaysIn) == - static_cast(DataClauseModifier::alwaysin) && - static_cast(OpenACCModifierKind::AlwaysOut) == - static_cast(DataClauseModifier::alwaysout) && - static_cast(OpenACCModifierKind::Capture) == - static_cast(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(modifiers); - return mlirModifiers; - } - template 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(OpenACCModifierKind::Zero) == + static_cast(DataClauseModifier::zero) && + static_cast(OpenACCModifierKind::Readonly) == + static_cast(DataClauseModifier::readonly) && + static_cast(OpenACCModifierKind::AlwaysIn) == + static_cast(DataClauseModifier::alwaysin) && + static_cast(OpenACCModifierKind::AlwaysOut) == + static_cast(DataClauseModifier::alwaysout) && + static_cast(OpenACCModifierKind::Capture) == + static_cast(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(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..e5cf70190b849 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 +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {name = "GlobalHSE1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// CHECK: acc.global_dtor @GlobalHSE1_acc_dtor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSE1 : !cir.ptr +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "GlobalHSE1"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr) {dataClause = #acc, 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 +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {name = "GlobalInt1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// CHECK: acc.global_dtor @GlobalInt1_acc_dtor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalInt1 : !cir.ptr +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "GlobalInt1"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr) {dataClause = #acc, 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> +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "GlobalHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr>) +// 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> +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "GlobalHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr>) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr>) {dataClause = #acc, 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 +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {modifiers = #acc, name = "NSHSE1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// CHECK: acc.global_dtor @{{.*}}NSHSE1{{.*}}_acc_dtor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "NSHSE1"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, 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 +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {modifiers = #acc, name = "NSInt1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// CHECK: acc.global_dtor @{{.*}}NSInt1{{.*}}_acc_dtor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "NSInt1"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, 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> +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {modifiers = #acc, name = "NSHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr>) +// 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> +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, modifiers = #acc, name = "NSHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr>) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr>) {dataClause = #acc, modifiers = #acc, 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 +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {name = "AnonNSHSE1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// CHECK: acc.global_dtor @{{.*}}AnonNSHSE1{{.*}}_acc_dtor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSE1{{.*}} : !cir.ptr +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "AnonNSHSE1"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr) {dataClause = #acc, 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 +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {name = "AnonNSInt1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr) +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } +// CHECK: acc.global_dtor @{{.*}}AnonNSInt1{{.*}}_acc_dtor { +// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSInt1{{.*}} : !cir.ptr +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "AnonNSInt1"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr) {dataClause = #acc, 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> +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "AnonNSHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr>) +// 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> +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "AnonNSHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr>) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr>) {dataClause = #acc, 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 +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {name = "StaticMemHSE1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr) +// 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 +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "StaticMemHSE1"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr) {dataClause = #acc, 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 +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {name = "StaticMemInt1"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr) +// 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 +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "StaticMemInt1"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr) {dataClause = #acc, 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> +// CHECK-NEXT: %[[CREATE:.*]] = acc.create varPtr(%[[GET_GLOBAL]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {name = "StaticMemHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_enter dataOperands(%[[CREATE]] : !cir.ptr>) +// 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> +// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr>) bounds(%[[BOUNDS]]) -> !cir.ptr> {dataClause = #acc, name = "StaticMemHSEArr[1:1]"} +// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr>) +// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr>) {dataClause = #acc, name = "StaticMemHSEArr[1:1]"} +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } void MemFunc1(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) { // CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr{{.*}}) 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)