diff --git a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
index aeb43f2c7bbed..0ee668a362b79 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp
@@ -247,6 +247,14 @@ class OpenACCGlobalDeclareClauseEmitter final
     this->VisitClauseList(clauses);
   }
 
+  void VisitCopyInClause(const OpenACCCopyInClause &clause) {
+    for (const Expr *var : clause.getVarList())
+      cgm.emitGlobalOpenACCDeclareDataOperands<mlir::acc::CopyinOp>(
+          var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
+          /*structured=*/true,
+          /*implicit=*/false, /*requiresDtor=*/true);
+  }
+
   void VisitCreateClause(const OpenACCCreateClause &clause) {
     for (const Expr *var : clause.getVarList())
       cgm.emitGlobalOpenACCDeclareDataOperands<mlir::acc::CreateOp>(
diff --git a/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp b/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp
index 1ed7a7d101adb..9f8e78a58a97d 100644
--- a/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp
@@ -5,14 +5,258 @@ struct HasSideEffects {
   ~HasSideEffects();
 };
 
-// TODO: OpenACC: Implement 'global', NS lowering.
+HasSideEffects GlobalHSE1;
+HasSideEffects GlobalHSEArr[5];
+int GlobalInt1;
+
+#pragma acc declare copyin(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: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "GlobalHSE1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !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_copyin>, 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_copyin>, 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: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "GlobalInt1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !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_copyin>, 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_copyin>, 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: %[[COPYIN:.*]] = acc.copyin 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(%[[COPYIN]] : !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_copyin>, 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_copyin>, name = "GlobalHSEArr[1:1]"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+
+namespace NS {
+
+HasSideEffects NSHSE1;
+HasSideEffects NSHSEArr[5];
+int NSInt1;
+
+#pragma acc declare copyin(alwaysin: 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: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {modifiers = #acc<data_clause_modifier alwaysin>, name = "NSHSE1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !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_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, 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_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, 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: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier alwaysin>, name = "NSInt1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !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_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, 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_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, 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: %[[COPYIN:.*]] = acc.copyin 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 alwaysin>, name = "NSHSEArr[1:1]"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !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_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, 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_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "NSHSEArr[1:1]"}
+// CHECK-NEXT: acc.terminator
+// CHECK-NEXT: }
+
+} // namespace NS
+
+namespace {
+
+HasSideEffects AnonNSHSE1;
+HasSideEffects AnonNSHSEArr[5];
+int AnonNSInt1;
+
+#pragma acc declare copyin(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: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "AnonNSHSE1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !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_copyin>, 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_copyin>, 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: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "AnonNSInt1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !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_copyin>, 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_copyin>, 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: %[[COPYIN:.*]] = acc.copyin 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(%[[COPYIN]] : !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_copyin>, 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_copyin>, 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 copyin(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: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "StaticMemHSE1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !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_copyin>, 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_copyin>, 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: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "StaticMemInt1"}
+// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !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_copyin>, 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_copyin>, 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: %[[COPYIN:.*]] = acc.copyin 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(%[[COPYIN]] : !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_copyin>, 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_copyin>, 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>{{.*}})