llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) <details> <summary>Changes</summary> This is the same as create/copyin, except it uses declare_device_resident for the entry op. --- Patch is 20.60 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/169507.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp (+9) - (modified) clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp (+248-4) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp index 0ee668a362b79..c1a1f8a83f5cd 100644 --- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp @@ -262,6 +262,15 @@ class OpenACCGlobalDeclareClauseEmitter final /*structured=*/true, /*implicit=*/false, /*requiresDtor=*/true); } + + void VisitDeviceResidentClause(const OpenACCDeviceResidentClause &clause) { + for (const Expr *var : clause.getVarList()) + cgm.emitGlobalOpenACCDeclareDataOperands< + mlir::acc::DeclareDeviceResidentOp>( + var, mlir::acc::DataClause::acc_declare_device_resident, {}, + /*structured=*/true, + /*implicit=*/false, /*requiresDtor=*/true); + } }; } // namespace diff --git a/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp b/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp index dbec4f22a1bb3..6640c2581d24d 100644 --- a/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp +++ b/clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp @@ -5,14 +5,258 @@ struct HasSideEffects { ~HasSideEffects(); }; -// TODO: OpenACC: Implement 'global', NS lowering. +HasSideEffects GlobalHSE1; +HasSideEffects GlobalHSEArr[5]; +int GlobalInt1; + +#pragma acc declare device_resident(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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, 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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, 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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, name = "GlobalHSEArr[1:1]"} +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } + +namespace NS { + +HasSideEffects NSHSE1; +HasSideEffects NSHSEArr[5]; +int NSInt1; + +#pragma acc declare device_resident(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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, 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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, 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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, name = "NSHSEArr[1:1]"} +// CHECK-NEXT: acc.terminator +// CHECK-NEXT: } + +} // namespace NS + +namespace { + +HasSideEffects AnonNSHSE1; +HasSideEffects AnonNSHSEArr[5]; +int AnonNSInt1; + +#pragma acc declare device_resident(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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, 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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, 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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, 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 device_resident(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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, 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.declare_device_resident 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_declare_device_resident>, 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_declare_device_resident>, 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.declare_device_resident 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_declare_device_resident>, name = "StaticMemHSEArr[1:1]"} +// CHECK-NEXT: acc.decla... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/169507 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
