From 97c3056371a458665272ae8d67898d12860028f9 Mon Sep 17 00:00:00 2001 From: erichkeane Date: Mon, 24 Nov 2025 10:23:30 -0800 Subject: [PATCH] [OpenACC][CIR] 'device_resident' clause lowering for local declare Just like the last handful of clauses, this is a pretty simple one, doing device_resident (Entry op: declare_device_resident, and exit: delete). This should be the last of the 'local' declare patches. --- clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp | 16 +- clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 12 ++ .../CodeGenOpenACC/declare-deviceresident.cpp | 199 ++++++++++++++++++ 3 files changed, 219 insertions(+), 8 deletions(-) create mode 100644 clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp index 9c1aeb87c8029..41a193e4d85c5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp @@ -19,12 +19,9 @@ using namespace clang::CIRGen; namespace { struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup { - SourceRange declareRange; mlir::acc::DeclareEnterOp enterOp; - OpenACCDeclareCleanup(SourceRange declareRange, - mlir::acc::DeclareEnterOp enterOp) - : declareRange(declareRange), enterOp(enterOp) {} + OpenACCDeclareCleanup(mlir::acc::DeclareEnterOp enterOp) : enterOp(enterOp) {} template void createOutOp(CIRGenFunction &cgf, InTy inOp) { @@ -78,8 +75,11 @@ struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup { createOutOp(cgf, create); break; } - } else if (auto create = val.getDefiningOp()) { - createOutOp(cgf, create); + } else if (auto present = val.getDefiningOp()) { + createOutOp(cgf, present); + } else if (auto dev_res = + val.getDefiningOp()) { + createOutOp(cgf, dev_res); } else if (val.getDefiningOp()) { // Link has no exit clauses, and shouldn't be copied. continue; @@ -87,7 +87,7 @@ struct OpenACCDeclareCleanup final : EHScopeStack::Cleanup { // DevicePtr has no exit clauses, and shouldn't be copied. continue; } else { - cgf.cgm.errorNYI(declareRange, "OpenACC local declare clause cleanup"); + llvm_unreachable("OpenACC local declare clause unexpected defining op"); continue; } exitOp.getDataClauseOperandsMutable().append(val); @@ -106,7 +106,7 @@ void CIRGenFunction::emitOpenACCDeclare(const OpenACCDeclareDecl &d) { d.clauses()); ehStack.pushCleanup(CleanupKind::NormalCleanup, - d.getSourceRange(), enterOp); + enterOp); } void CIRGenFunction::emitOpenACCRoutine(const OpenACCRoutineDecl &d) { diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index a23ec93ab1d75..60a089fe0e936 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -1135,6 +1135,18 @@ class OpenACCClauseCIREmitter final llvm_unreachable("Unknown construct kind in VisitReductionClause"); } } + + void VisitDeviceResidentClause(const OpenACCDeviceResidentClause &clause) { + if constexpr (isOneOfTypes) { + for (const Expr *var : clause.getVarList()) + addDataOperand( + var, mlir::acc::DataClause::acc_declare_device_resident, {}, + /*structured=*/true, + /*implicit=*/false); + } else { + llvm_unreachable("Unknown construct kind in VisitDeviceResidentClause"); + } + } }; template diff --git a/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp b/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp new file mode 100644 index 0000000000000..dbec4f22a1bb3 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.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 device_resident(ArgHSE, ArgInt, LocalHSE, LocalInt, ArgHSEPtr[1:1], LocalHSEArr[1:1]) + // CHECK: %[[ARG_HSE_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {name = "ArgHSE"} + // CHECK-NEXT: %[[ARG_INT_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {name = "ArgInt"} + // CHECK-NEXT: %[[LOC_HSE_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {name = "LocalHSE"} + // CHECK-NEXT: %[[LOC_INT_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {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_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND1]]) -> !cir.ptr> {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_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND2]]) -> !cir.ptr> {name = "LocalHSEArr[1:1]"} + // CHECK-NEXT: %[[ENTER:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr, !cir.ptr, !cir.ptr, !cir.ptr, !cir.ptr>, !cir.ptr>) + // + // CHECK-NEXT: acc.declare_exit token(%[[ENTER]]) dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr, !cir.ptr, !cir.ptr, !cir.ptr, !cir.ptr>, !cir.ptr>) + // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "ArgHSE"} + // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "ArgInt"} + // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "LocalHSE"} + // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "LocalInt"} + // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr>) bounds(%[[BOUND1]]) {dataClause = #acc, name = "ArgHSEPtr[1:1]"} + // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr>) bounds(%[[BOUND2]]) {dataClause = #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 device_resident(ArgHSE, ArgInt, ArgHSEPtr[1:1]) + // CHECK: %[[ARG_HSE_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {name = "ArgHSE"} + // CHECK-NEXT: %[[ARG_INT_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {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_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND1]]) -> !cir.ptr> {name = "ArgHSEPtr[1:1]"} + // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr, !cir.ptr, !cir.ptr>) + +#pragma acc declare device_resident(LocalHSE, LocalInt, LocalHSEArr[1:1]) + // CHECK-NEXT: %[[LOC_HSE_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {name = "LocalHSE"} + // CHECK-NEXT: %[[LOC_INT_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {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_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND2]]) -> !cir.ptr> {name = "LocalHSEArr[1:1]"} + // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr, !cir.ptr, !cir.ptr>) + + // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr, !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "LocalHSE"} + // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "LocalInt"} + // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr>) bounds(%[[BOUND2]]) {dataClause = #acc, name = "LocalHSEArr[1:1]"} + // + // CHECK-NEXT: acc.declare_exit token(%[[ENTER1]]) dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr, !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "ArgHSE"} + // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "ArgInt"} + // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr>) bounds(%[[BOUND1]]) {dataClause = #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 device_resident(ArgHSE, ArgInt, ArgHSEPtr[1:1]) + // CHECK: %[[ARG_HSE_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[ARG_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {name = "ArgHSE"} + // CHECK-NEXT: %[[ARG_INT_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[ARG_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {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_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[ARG_HSE_PTR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND1]]) -> !cir.ptr> {name = "ArgHSEPtr[1:1]"} + // CHECK-NEXT: %[[ENTER1:.*]] = acc.declare_enter dataOperands(%[[ARG_HSE_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr, !cir.ptr, !cir.ptr>) + { + // CHECK-NEXT: cir.scope { +#pragma acc declare device_resident(LocalHSE, LocalInt, LocalHSEArr[1:1]) + // CHECK-NEXT: %[[LOC_HSE_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ALLOCA]] : !cir.ptr) -> !cir.ptr {name = "LocalHSE"} + // CHECK-NEXT: %[[LOC_INT_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_INT_ALLOCA]] : !cir.ptr) -> !cir.ptr {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_DEV_RES:.*]] = acc.declare_device_resident varPtr(%[[LOC_HSE_ARR_ALLOCA]] : !cir.ptr>) bounds(%[[BOUND2]]) -> !cir.ptr> {name = "LocalHSEArr[1:1]"} + // CHECK-NEXT: %[[ENTER2:.*]] = acc.declare_enter dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr, !cir.ptr, !cir.ptr>) + + do_thing(); + // CHECK-NEXT: cir.call @do_thing + // CHECK-NEXT: acc.declare_exit token(%[[ENTER2]]) dataOperands(%[[LOC_HSE_DEV_RES]], %[[LOC_INT_DEV_RES]], %[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr, !cir.ptr, !cir.ptr>) + // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "LocalHSE"} + // CHECK-NEXT: acc.delete accPtr(%[[LOC_INT_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "LocalInt"} + // CHECK-NEXT: acc.delete accPtr(%[[LOC_HSE_ARR_DEV_RES]] : !cir.ptr>) bounds(%[[BOUND2]]) {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_DEV_RES]], %[[ARG_INT_DEV_RES]], %[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr, !cir.ptr, !cir.ptr>) + + // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "ArgHSE"} + // CHECK-NEXT: acc.delete accPtr(%[[ARG_INT_DEV_RES]] : !cir.ptr) {dataClause = #acc, name = "ArgInt"} + // CHECK-NEXT: acc.delete accPtr(%[[ARG_HSE_PTR_DEV_RES]] : !cir.ptr>) bounds(%[[BOUND1]]) {dataClause = #acc, name = "ArgHSEPtr[1:1]"} +} +