Skip to content

Commit be2dfce

Browse files
erichkeaneerichkeane
authored
[OpenACC][CIR] Global declare 'copyin' clause lowering (#169498)
JUST like the 'create' clause, except the entry op is copyin instead of create. Most of this is the test.
1 parent eb5297e commit be2dfce

File tree

2 files changed

+256
-4
lines changed

2 files changed

+256
-4
lines changed

‎clang/lib/CIR/CodeGen/CIRGenDeclOpenACC.cpp‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -247,6 +247,14 @@ class OpenACCGlobalDeclareClauseEmitter final
247247
this->VisitClauseList(clauses);
248248
}
249249

250+
void VisitCopyInClause(const OpenACCCopyInClause &clause) {
251+
for (const Expr *var : clause.getVarList())
252+
cgm.emitGlobalOpenACCDeclareDataOperands<mlir::acc::CopyinOp>(
253+
var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
254+
/*structured=*/true,
255+
/*implicit=*/false, /*requiresDtor=*/true);
256+
}
257+
250258
void VisitCreateClause(const OpenACCCreateClause &clause) {
251259
for (const Expr *var : clause.getVarList())
252260
cgm.emitGlobalOpenACCDeclareDataOperands<mlir::acc::CreateOp>(

‎clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp‎

Lines changed: 248 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5,14 +5,258 @@ struct HasSideEffects {
55
~HasSideEffects();
66
};
77

8-
// TODO: OpenACC: Implement 'global', NS lowering.
8+
HasSideEffects GlobalHSE1;
9+
HasSideEffects GlobalHSEArr[5];
10+
int GlobalInt1;
11+
12+
#pragma acc declare copyin(GlobalHSE1, GlobalInt1, GlobalHSEArr[1:1])
13+
// CHECK: acc.global_ctor @GlobalHSE1_acc_ctor {
14+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSE1 : !cir.ptr<!rec_HasSideEffects>
15+
// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "GlobalHSE1"}
16+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!rec_HasSideEffects>)
17+
// CHECK-NEXT: acc.terminator
18+
// CHECK-NEXT: }
19+
// CHECK: acc.global_dtor @GlobalHSE1_acc_dtor {
20+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSE1 : !cir.ptr<!rec_HasSideEffects>
21+
// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copyin>, name = "GlobalHSE1"}
22+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
23+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, name = "GlobalHSE1"}
24+
// CHECK-NEXT: acc.terminator
25+
// CHECK-NEXT: }
26+
//
27+
// CHECK: acc.global_ctor @GlobalInt1_acc_ctor {
28+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalInt1 : !cir.ptr<!s32i>
29+
// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "GlobalInt1"}
30+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!s32i>)
31+
// CHECK-NEXT: acc.terminator
32+
// CHECK-NEXT: }
33+
// CHECK: acc.global_dtor @GlobalInt1_acc_dtor {
34+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalInt1 : !cir.ptr<!s32i>
35+
// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyin>, name = "GlobalInt1"}
36+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
37+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "GlobalInt1"}
38+
// CHECK-NEXT: acc.terminator
39+
// CHECK-NEXT: }
40+
//
41+
// CHECK: acc.global_ctor @GlobalHSEArr_acc_ctor {
42+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
43+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
44+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
45+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
46+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
47+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
48+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
49+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
50+
// 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]"}
51+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
52+
// CHECK-NEXT: acc.terminator
53+
// CHECK-NEXT: }
54+
// CHECK: acc.global_dtor @GlobalHSEArr_acc_dtor {
55+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
56+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
57+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
58+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
59+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
60+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
61+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
62+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @GlobalHSEArr : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
63+
// 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]"}
64+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
65+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_copyin>, name = "GlobalHSEArr[1:1]"}
66+
// CHECK-NEXT: acc.terminator
67+
// CHECK-NEXT: }
68+
69+
namespace NS {
70+
71+
HasSideEffects NSHSE1;
72+
HasSideEffects NSHSEArr[5];
73+
int NSInt1;
74+
75+
#pragma acc declare copyin(alwaysin: NSHSE1, NSInt1, NSHSEArr[1:1])
76+
// CHECK: acc.global_ctor @{{.*}}NSHSE1{{.*}}_acc_ctor {
77+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
78+
// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {modifiers = #acc<data_clause_modifier alwaysin>, name = "NSHSE1"}
79+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!rec_HasSideEffects>)
80+
// CHECK-NEXT: acc.terminator
81+
// CHECK-NEXT: }
82+
// CHECK: acc.global_dtor @{{.*}}NSHSE1{{.*}}_acc_dtor {
83+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
84+
// 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"}
85+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
86+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "NSHSE1"}
87+
// CHECK-NEXT: acc.terminator
88+
// CHECK-NEXT: }
89+
//
90+
// CHECK: acc.global_ctor @{{.*}}NSInt1{{.*}}_acc_ctor {
91+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
92+
// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier alwaysin>, name = "NSInt1"}
93+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!s32i>)
94+
// CHECK-NEXT: acc.terminator
95+
// CHECK-NEXT: }
96+
// CHECK: acc.global_dtor @{{.*}}NSInt1{{.*}}_acc_dtor {
97+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSInt1{{.*}} : !cir.ptr<!s32i>
98+
// 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"}
99+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
100+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier alwaysin>, name = "NSInt1"}
101+
// CHECK-NEXT: acc.terminator
102+
// CHECK-NEXT: }
103+
//
104+
// CHECK: acc.global_ctor @{{.*}}NSHSEArr{{.*}}_acc_ctor {
105+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
106+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
107+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
108+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
109+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
110+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
111+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
112+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
113+
// 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]"}
114+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
115+
// CHECK-NEXT: acc.terminator
116+
// CHECK-NEXT: }
117+
// CHECK: acc.global_dtor @{{.*}}NSHSEArr{{.*}}_acc_dtor {
118+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
119+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
120+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
121+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
122+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
123+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
124+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
125+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}NSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
126+
// 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]"}
127+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
128+
// 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]"}
129+
// CHECK-NEXT: acc.terminator
130+
// CHECK-NEXT: }
131+
132+
} // namespace NS
133+
134+
namespace {
135+
136+
HasSideEffects AnonNSHSE1;
137+
HasSideEffects AnonNSHSEArr[5];
138+
int AnonNSInt1;
139+
140+
#pragma acc declare copyin(AnonNSHSE1, AnonNSInt1, AnonNSHSEArr[1:1])
141+
// CHECK: acc.global_ctor @{{.*}}AnonNSHSE1{{.*}}_acc_ctor {
142+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
143+
// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "AnonNSHSE1"}
144+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!rec_HasSideEffects>)
145+
// CHECK-NEXT: acc.terminator
146+
// CHECK-NEXT: }
147+
// CHECK: acc.global_dtor @{{.*}}AnonNSHSE1{{.*}}_acc_dtor {
148+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
149+
// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copyin>, name = "AnonNSHSE1"}
150+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
151+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, name = "AnonNSHSE1"}
152+
// CHECK-NEXT: acc.terminator
153+
// CHECK-NEXT: }
154+
//
155+
// CHECK: acc.global_ctor @{{.*}}AnonNSInt1{{.*}}_acc_ctor {
156+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSInt1{{.*}} : !cir.ptr<!s32i>
157+
// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "AnonNSInt1"}
158+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!s32i>)
159+
// CHECK-NEXT: acc.terminator
160+
// CHECK-NEXT: }
161+
// CHECK: acc.global_dtor @{{.*}}AnonNSInt1{{.*}}_acc_dtor {
162+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSInt1{{.*}} : !cir.ptr<!s32i>
163+
// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyin>, name = "AnonNSInt1"}
164+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
165+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "AnonNSInt1"}
166+
// CHECK-NEXT: acc.terminator
167+
// CHECK-NEXT: }
168+
//
169+
// CHECK: acc.global_ctor @{{.*}}AnonNSHSEArr{{.*}}_acc_ctor {
170+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
171+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
172+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
173+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
174+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
175+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
176+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
177+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
178+
// 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]"}
179+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
180+
// CHECK-NEXT: acc.terminator
181+
// CHECK-NEXT: }
182+
// CHECK: acc.global_dtor @{{.*}}AnonNSHSEArr{{.*}}_acc_dtor {
183+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
184+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
185+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
186+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
187+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
188+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
189+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
190+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}AnonNSHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
191+
// 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]"}
192+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
193+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_copyin>, name = "AnonNSHSEArr[1:1]"}
194+
// CHECK-NEXT: acc.terminator
195+
// CHECK-NEXT: }
196+
197+
} // namespace NS
9198

10199
struct Struct {
11-
static const HasSideEffects StaticMemHSE;
200+
static const HasSideEffects StaticMemHSE1;
12201
static const HasSideEffects StaticMemHSEArr[5];
13-
static const int StaticMemInt;
202+
static const int StaticMemInt1;
14203

15-
// TODO: OpenACC: Implement static-local lowering.
204+
#pragma acc declare copyin(StaticMemHSE1, StaticMemInt1, StaticMemHSEArr[1:1])
205+
// CHECK: acc.global_ctor @{{.*}}Struct{{.*}}StaticMemHSE1{{.*}}_acc_ctor {
206+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}{{.*}}Struct{{.*}}StaticMemHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
207+
// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {name = "StaticMemHSE1"}
208+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!rec_HasSideEffects>)
209+
// CHECK-NEXT: acc.terminator
210+
// CHECK-NEXT: }
211+
// CHECK: acc.global_dtor @{{.*}}Struct{{.*}}StaticMemHSE1{{.*}}_acc_dtor {
212+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}Struct{{.*}}StaticMemHSE1{{.*}} : !cir.ptr<!rec_HasSideEffects>
213+
// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!rec_HasSideEffects>) -> !cir.ptr<!rec_HasSideEffects> {dataClause = #acc<data_clause acc_copyin>, name = "StaticMemHSE1"}
214+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>)
215+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!rec_HasSideEffects>) {dataClause = #acc<data_clause acc_copyin>, name = "StaticMemHSE1"}
216+
// CHECK-NEXT: acc.terminator
217+
// CHECK-NEXT: }
218+
//
219+
// CHECK: acc.global_ctor @{{.*}}Struct{{.*}}StaticMemInt1{{.*}}_acc_ctor {
220+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}Struct{{.*}}StaticMemInt1{{.*}} : !cir.ptr<!s32i>
221+
// CHECK-NEXT: %[[COPYIN:.*]] = acc.copyin varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "StaticMemInt1"}
222+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!s32i>)
223+
// CHECK-NEXT: acc.terminator
224+
// CHECK-NEXT: }
225+
// CHECK: acc.global_dtor @{{.*}}Struct{{.*}}StaticMemInt1{{.*}}_acc_dtor {
226+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}Struct{{.*}}StaticMemInt1{{.*}} : !cir.ptr<!s32i>
227+
// CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[GET_GLOBAL]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyin>, name = "StaticMemInt1"}
228+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
229+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "StaticMemInt1"}
230+
// CHECK-NEXT: acc.terminator
231+
// CHECK-NEXT: }
232+
//
233+
// CHECK: acc.global_ctor @{{.*}}Struct{{.*}}StaticMemHSEArr{{.*}}_acc_ctor {
234+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
235+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
236+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
237+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
238+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
239+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
240+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
241+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}Struct{{.*}}StaticMemHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
242+
// 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]"}
243+
// CHECK-NEXT: acc.declare_enter dataOperands(%[[COPYIN]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
244+
// CHECK-NEXT: acc.terminator
245+
// CHECK-NEXT: }
246+
// CHECK: acc.global_dtor @{{.*}}Struct{{.*}}StaticMemHSEArr{{.*}}_acc_dtor {
247+
// CHECK-NEXT: %[[LB:.*]] = cir.const #cir.int<1> : !s32i
248+
// CHECK-NEXT: %[[LB_CAST:.*]] = builtin.unrealized_conversion_cast %[[LB]]
249+
// CHECK-NEXT: %[[UB:.*]] = cir.const #cir.int<1> : !s32i
250+
// CHECK-NEXT: %[[UB_CAST:.*]] = builtin.unrealized_conversion_cast %[[UB]]
251+
// CHECK-NEXT: %[[IDX:.*]] = arith.constant 0 : i64
252+
// CHECK-NEXT: %[[STRIDE:.*]] = arith.constant 1 : i64
253+
// CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[LB_CAST]] : si32) extent(%[[UB_CAST]] : si32) stride(%[[STRIDE]] : i64) startIdx(%[[IDX]] : i64)
254+
// CHECK-NEXT: %[[GET_GLOBAL:.*]] = cir.get_global @{{.*}}Struct{{.*}}StaticMemHSEArr{{.*}} : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>
255+
// 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]"}
256+
// CHECK-NEXT: acc.declare_exit dataOperands(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>)
257+
// CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!cir.array<!rec_HasSideEffects x 5>>) {dataClause = #acc<data_clause acc_copyin>, name = "StaticMemHSEArr[1:1]"}
258+
// CHECK-NEXT: acc.terminator
259+
// CHECK-NEXT: }
16260

17261
void MemFunc1(HasSideEffects ArgHSE, int ArgInt, HasSideEffects *ArgHSEPtr) {
18262
// CHECK: cir.func {{.*}}MemFunc1{{.*}}(%{{.*}}: !cir.ptr<!rec_Struct>{{.*}}, %[[ARG_HSE:.*]]: !rec_HasSideEffects{{.*}}, %[[ARG_INT:.*]]: !s32i {{.*}}, %[[ARG_HSE_PTR:.*]]: !cir.ptr<!rec_HasSideEffects>{{.*}})

0 commit comments

Comments
 (0)