diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp index 551027bb1c8eb..581a6ca81e2c4 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp @@ -19,18 +19,52 @@ using namespace clang::CIRGen; namespace { struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup { + SourceRange declareRange; mlir::acc::DeclareEnterOp enterOp; - OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp) : enterOp(enterOp) {} + OpenACCDeclareCleanup(SourceRange declareRange, + mlir::acc::DeclareEnterOp enterOp) + : declareRange(declareRange), enterOp(enterOp) {} + + template + void createOutOp(CIRGenFunction &cgf, InTy inOp) { + auto outOp = + OutTy::create(cgf.getBuilder(), inOp.getLoc(), inOp, inOp.getVarPtr(), + inOp.getStructured(), inOp.getImplicit(), + llvm::Twine(inOp.getNameAttr()), inOp.getBounds()); + outOp.setDataClause(inOp.getDataClause()); + outOp.setModifiers(inOp.getModifiers()); + } void emit(CIRGenFunction &cgf) override { - mlir::acc::DeclareExitOp::create(cgf.getBuilder(), enterOp.getLoc(), - enterOp, {}); + auto exitOp = mlir::acc::DeclareExitOp::create( + cgf.getBuilder(), enterOp.getLoc(), enterOp, {}); - // TODO(OpenACC): Some clauses require that we add info about them to the - // DeclareExitOp. However, we don't have any of those implemented yet, so - // we should add infrastructure here to do that once we have one - // implemented. + // Some data clauses need to be referenced in 'exit', AND need to have an + // operation after the exit. Copy these from the enter operation. + for (mlir::Value val : enterOp.getDataClauseOperands()) { + if (auto copyin = val.getDefiningOp()) { + switch (copyin.getDataClause()) { + default: + cgf.cgm.errorNYI(declareRange, + "OpenACC local declare clause copyin cleanup"); + break; + case mlir::acc::DataClause::acc_copy: + createOutOp(cgf, copyin); + break; + } + } else if (val.getDefiningOp()) { + // Link has no exit clauses, and shouldn't be copied. + continue; + } else if (val.getDefiningOp()) { + // DevicePtr has no exit clauses, and shouldn't be copied. + continue; + } else { + cgf.cgm.errorNYI(declareRange, "OpenACC local declare clause cleanup"); + continue; + } + exitOp.getDataClauseOperandsMutable().append(val); + } } }; } // namespace @@ -45,7 +79,7 @@ void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) { d.clauses()); ehStack.pushCleanup(CleanupKind::NormalCleanup, - enterOp); + d.getSourceRange(), enterOp); } void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) { diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index c5c6bcd0153a4..621af2344209f 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -800,12 +800,16 @@ class OpenACCClauseCIREmitter final var, mlir::acc::DataClause::acc_copy, clause.getModifierList(), /*structured=*/true, /*implicit=*/false); + } else if constexpr (isOneOfTypes) { + for (const Expr *var : clause.getVarList()) + addDataOperand( + var, mlir::acc::DataClause::acc_copy, clause.getModifierList(), + /*structured=*/true, + /*implicit=*/false); } else if constexpr (isCombinedType) { applyToComputeOp(clause); } else { - // TODO: When we've implemented this for everything, switch this to an - // unreachable. declare construct remains. - return clauseNotImplemented(clause); + llvm_unreachable("Unknown construct kind in VisitCopyClause"); } } diff --git a/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp b/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp new file mode 100644 index 0000000000000..a8a9115a21b29 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/declare-copy.cpp @@ -0,0 +1,199 @@ +// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s + +struct HasSideEffects { + HasSideEffects(); + ~HasSideEffects(); +}; + +// TODO: OpenACC: Implement 'global', NS lowering. + +struct Struct { + static const HasSideEffects StaticMemHSE; + static const HasSideEffects StaticMemHSEArr[5]; + static const int StaticMemInt; + + // TODO: OpenACC: Implement static-local lowering. + + void MemFunc1(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) { + // CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr{{.*}}) + // CHECK-NEXT: cir.alloca{{.*}}["this" + // CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !rec_HasSideEffects{{.*}}["ArgHSE" + // CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !s32i{{.*}}["ArgInt + // CHECK-NEXT: %[[ARG_HSE_PTR_ALLOCA:.*]] = cir.alloca !cir.ptr{{.*}}["ArgHSEPtr" + // CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !rec_HasSideEffects{{.*}}["LocalHSE + // CHECK-NEXT: %[[LOC_HSE_ARR_ALLOCA:.*]] = cir.alloca !cir.array{{.*}}["LocalHSEArr + // CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !s32i{{.*}}["LocalInt + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.load + + HasSideEffects LocalHSE; + // CHECK-NEXT: cir.call{{.*}} : (!cir.ptr) -> () + HasSideEffects LocalHSEArr[5]; + int LocalInt; + +#pragma acc declare copy(always:ArgHSE, ArgInt, LocalHSE, LocalInt, ArgHSEPtr[1:1], LocalHSEArr[1:1]) + // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "ArgHSE"} + // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "ArgInt"} + // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "LocalHSE"} + // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "LocalInt"} + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64) + // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND1]]) -> !cir.ptr> {dataClause = #acc, modifiers = #acc, name = "ArgHSEPtr[1:1]"} + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64) + // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND2]]) -> !cir.ptr> {dataClause = #acc, modifiers = #acc, name = "LocalHSEArr[1:1]"} + // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr, !cir.ptr, !cir.ptr, !cir.ptr, !cir.ptr>, !cir.ptr>) + // + // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr, !cir.ptr, !cir.ptr, !cir.ptr, !cir.ptr>, !cir.ptr>) + // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, name = "ArgHSE"} + // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, name = "ArgInt"} + // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, name = "LocalHSE"} + // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, name = "LocalInt"} + // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr>) {dataClause = #acc, modifiers = #acc, name = "ArgHSEPtr[1:1]"} + // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr>) {dataClause = #acc, modifiers = #acc, name = "LocalHSEArr[1:1]"} + } + void MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr); +}; + +void use() { + Struct s; + s.MemFunc1(HasSideEffects{}, 0, nullptr); +} + +void Struct::MemFunc2(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) { + // CHECK: cir.func {{.*}}MemFunc2{{.*}}(%{{.*}}: !cir.ptr{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr{{.*}}) + // CHECK-NEXT: cir.alloca{{.*}}["this" + // CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !rec_HasSideEffects{{.*}}["ArgHSE" + // CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !s32i{{.*}}["ArgInt + // CHECK-NEXT: %[[ARG_HSE_PTR_ALLOCA:.*]] = cir.alloca !cir.ptr{{.*}}["ArgHSEPtr" + // CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !rec_HasSideEffects{{.*}}["LocalHSE + // CHECK-NEXT: %[[LOC_HSE_ARR_ALLOCA:.*]] = cir.alloca !cir.array{{.*}}["LocalHSEArr + // CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !s32i{{.*}}["LocalInt + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.load + HasSideEffects LocalHSE; + // CHECK-NEXT: cir.call{{.*}} : (!cir.ptr) -> () + HasSideEffects LocalHSEArr[5]; + // CHECK: do { + // CHECK: } while { + // CHECK: } + int LocalInt; +#pragma acc declare copy(alwaysin:ArgHSE, ArgInt, ArgHSEPtr[1:1]) + // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "ArgHSE"} + // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "ArgInt"} + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64) + // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND1]]) -> !cir.ptr> {dataClause = #acc, modifiers = #acc, name = "ArgHSEPtr[1:1]"} + // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr, !cir.ptr, !cir.ptr>) + +#pragma acc declare copy(alwaysout:LocalHSE, LocalInt, LocalHSEArr[1:1]) + // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "LocalHSE"} + // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "LocalInt"} + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64) + // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND2]]) -> !cir.ptr> {dataClause = #acc, modifiers = #acc, name = "LocalHSEArr[1:1]"} + // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr, !cir.ptr, !cir.ptr>) + + // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr, !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, name = "LocalHSE"} + // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, name = "LocalInt"} + // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr>) {dataClause = #acc, modifiers = #acc, name = "LocalHSEArr[1:1]"} + // + // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr, !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, name = "ArgHSE"} + // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, name = "ArgInt"} + // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr>) {dataClause = #acc, modifiers = #acc, name = "ArgHSEPtr[1:1]"} +} + +extern "C" void do_thing(); + +extern "C" void NormalFunc(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) { + // CHECK: cir.func {{.*}}NormalFunc(%[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr{{.*}}) + // CHECK-NEXT: %[[ARG_HSE_ALLOCA:.*]] = cir.alloca !rec_HasSideEffects{{.*}}["ArgHSE" + // CHECK-NEXT: %[[ARG_INT_ALLOCA:.*]] = cir.alloca !s32i{{.*}}["ArgInt + // CHECK-NEXT: %[[ARG_HSE_PTR_ALLOCA:.*]] = cir.alloca !cir.ptr{{.*}}["ArgHSEPtr" + // CHECK-NEXT: %[[LOC_HSE_ALLOCA:.*]] = cir.alloca !rec_HasSideEffects{{.*}}["LocalHSE + // CHECK-NEXT: %[[LOC_HSE_ARR_ALLOCA:.*]] = cir.alloca !cir.array{{.*}}["LocalHSEArr + // CHECK-NEXT: %[[LOC_INT_ALLOCA:.*]] = cir.alloca !s32i{{.*}}["LocalInt + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.store + // CHECK-NEXT: cir.store + HasSideEffects LocalHSE; + // CHECK-NEXT: cir.call{{.*}} : (!cir.ptr) -> () + HasSideEffects LocalHSEArr[5]; + // CHECK: do { + // CHECK: } while { + // CHECK: } + int LocalInt; +#pragma acc declare copy(capture:ArgHSE, ArgInt, ArgHSEPtr[1:1]) + // CHECK: %[[ARG_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "ArgHSE"} + // CHECK-NEXT: %[[ARG_INT_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, modifiers = #acc, name = "ArgInt"} + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUND1:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64) + // CHECK-NEXT: %[[ARG_HSE_PTR_COPYIN:.*]] = acc.copyin varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND1]]) -> !cir.ptr> {dataClause = #acc, modifiers = #acc, name = "ArgHSEPtr[1:1]"} + // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr, !cir.ptr, !cir.ptr>) + { + // CHECK-NEXT: cir.scope { +#pragma acc declare copy(LocalHSE, LocalInt, LocalHSEArr[1:1]) + // CHECK-NEXT: %[[LOC_HSE_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "LocalHSE"} + // CHECK-NEXT: %[[LOC_INT_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {dataClause = #acc, name = "LocalInt"} + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[LB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i + // CHECK-NEXT: %[[UB:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64 + // CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64 + // CHECK-NEXT: %[[BOUND2:.*]] = acc.bounds lowerbound(%[[LB]] : si32) extent(%[[UB]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64) + // CHECK-NEXT: %[[LOC_HSE_ARR_COPYIN:.*]] = acc.copyin varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND2]]) -> !cir.ptr> {dataClause = #acc, name = "LocalHSEArr[1:1]"} + // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr, !cir.ptr, !cir.ptr>) + + do_thing(); + // CHECK-NEXT: cir.call @do_thing + // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_COPYIN]], %[[LOC_INT_COPYIN]], %[[LOC_HSE_ARR_COPYIN]] : !cir.ptr, !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_COPYIN]] : !cir.ptr) to varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr) {dataClause = #acc, name = "LocalHSE"} + // CHECK-NEXT: acc.copyout accPtr(%[[LOC_INT_COPYIN]] : !cir.ptr) to varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr) {dataClause = #acc, name = "LocalInt"} + // CHECK-NEXT: acc.copyout accPtr(%[[LOC_HSE_ARR_COPYIN]] : !cir.ptr>) bounds(%[[BOUND2]]) to varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr>) {dataClause = #acc, name = "LocalHSEArr[1:1]"} + } + // CHECK-NEXT: } + + // Make sure that cleanup gets put in the right scope. + do_thing(); + // CHECK-NEXT: cir.call @do_thing + // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_COPYIN]], %[[ARG_INT_COPYIN]], %[[ARG_HSE_PTR_COPYIN]] : !cir.ptr, !cir.ptr, !cir.ptr>) + + // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_COPYIN]] : !cir.ptr) to varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, name = "ArgHSE"} + // CHECK-NEXT: acc.copyout accPtr(%[[ARG_INT_COPYIN]] : !cir.ptr) to varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr) {dataClause = #acc, modifiers = #acc, name = "ArgInt"} + // CHECK-NEXT: acc.copyout accPtr(%[[ARG_HSE_PTR_COPYIN]] : !cir.ptr>) bounds(%[[BOUND1]]) to varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr>) {dataClause = #acc, modifiers = #acc, name = "ArgHSEPtr[1:1]"} +} +