Skip to content

Commit b76bc18

Browse files
erichkeaneerichkeane
committed
[OpenACC][CIR] Add copy/etc clause lowering for 'data'.
These work exactly the same way they do for compute constructs, so this implements them identically and adds tests. The list of legal modifiers is different, but that is all handled in Sema.
1 parent e933cfc commit b76bc18

File tree

2 files changed

+198
-8
lines changed

2 files changed

+198
-8
lines changed

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

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -843,7 +843,7 @@ class OpenACCClauseCIREmitter final
843843

844844
void VisitCopyClause(const OpenACCCopyClause &clause) {
845845
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
846-
mlir::acc::KernelsOp>) {
846+
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
847847
for (const Expr *var : clause.getVarList())
848848
addDataOperand<mlir::acc::CopyinOp, mlir::acc::CopyoutOp>(
849849
var, mlir::acc::DataClause::acc_copy, clause.getModifierList(),
@@ -853,14 +853,14 @@ class OpenACCClauseCIREmitter final
853853
applyToComputeOp(clause);
854854
} else {
855855
// TODO: When we've implemented this for everything, switch this to an
856-
// unreachable. data, declare, combined constructs remain.
856+
// unreachable. declare construct remains.
857857
return clauseNotImplemented(clause);
858858
}
859859
}
860860

861861
void VisitCopyInClause(const OpenACCCopyInClause &clause) {
862862
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
863-
mlir::acc::KernelsOp>) {
863+
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
864864
for (const Expr *var : clause.getVarList())
865865
addDataOperand<mlir::acc::CopyinOp, mlir::acc::DeleteOp>(
866866
var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
@@ -870,14 +870,14 @@ class OpenACCClauseCIREmitter final
870870
applyToComputeOp(clause);
871871
} else {
872872
// TODO: When we've implemented this for everything, switch this to an
873-
// unreachable. data, declare, combined constructs remain.
873+
// unreachable. enter-data, declare constructs remain.
874874
return clauseNotImplemented(clause);
875875
}
876876
}
877877

878878
void VisitCopyOutClause(const OpenACCCopyOutClause &clause) {
879879
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
880-
mlir::acc::KernelsOp>) {
880+
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
881881
for (const Expr *var : clause.getVarList())
882882
addDataOperand<mlir::acc::CreateOp, mlir::acc::CopyoutOp>(
883883
var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
@@ -887,14 +887,14 @@ class OpenACCClauseCIREmitter final
887887
applyToComputeOp(clause);
888888
} else {
889889
// TODO: When we've implemented this for everything, switch this to an
890-
// unreachable. data, declare, combined constructs remain.
890+
// unreachable. exit data, declare constructs remain.
891891
return clauseNotImplemented(clause);
892892
}
893893
}
894894

895895
void VisitCreateClause(const OpenACCCreateClause &clause) {
896896
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
897-
mlir::acc::KernelsOp>) {
897+
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
898898
for (const Expr *var : clause.getVarList())
899899
addDataOperand<mlir::acc::CreateOp, mlir::acc::DeleteOp>(
900900
var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
@@ -904,7 +904,7 @@ class OpenACCClauseCIREmitter final
904904
applyToComputeOp(clause);
905905
} else {
906906
// TODO: When we've implemented this for everything, switch this to an
907-
// unreachable. data, declare, combined constructs remain.
907+
// unreachable. enter-data, declare constructs remain.
908908
return clauseNotImplemented(clause);
909909
}
910910
}

‎clang/test/CIR/CodeGenOpenACC/data-copy-copyin-copyout-create.c

Lines changed: 190 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,190 @@
1+
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
2+
3+
void acc_data(int parmVar) {
4+
// CHECK: cir.func{{.*}} @acc_data(%[[ARG:.*]]: !s32i{{.*}}) {
5+
// CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar", init]
6+
int localVar1;
7+
// CHECK-NEXT: %[[LV1:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["localVar1"]
8+
// CHECK-NEXT: cir.store %[[ARG]], %[[PARM]]
9+
10+
#pragma acc data copy(parmVar)
11+
;
12+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "parmVar"}
13+
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
14+
// CHECK-NEXT: acc.terminator
15+
// CHECK-NEXT: } loc
16+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"}
17+
#pragma acc data copy(parmVar) copy(localVar1)
18+
;
19+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "parmVar"}
20+
// CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"}
21+
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
22+
// CHECK-NEXT: acc.terminator
23+
// CHECK-NEXT: } loc
24+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"}
25+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"}
26+
#pragma acc data copy(parmVar, localVar1)
27+
;
28+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "parmVar"}
29+
// CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"}
30+
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
31+
// CHECK-NEXT: acc.terminator
32+
// CHECK-NEXT: } loc
33+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "localVar1"}
34+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, name = "parmVar"}
35+
36+
#pragma acc data copy(capture, always, alwaysin, alwaysout: parmVar)
37+
;
38+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always,capture>, name = "parmVar"}
39+
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
40+
// CHECK-NEXT: acc.terminator
41+
// CHECK-NEXT: } loc
42+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always,capture>, name = "parmVar"}
43+
44+
#pragma acc data copy(capture: parmVar) copy(always, alwaysin, alwaysout: parmVar)
45+
;
46+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
47+
// CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "parmVar"}
48+
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
49+
// CHECK-NEXT: acc.terminator
50+
// CHECK-NEXT: } loc
51+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier always>, name = "parmVar"}
52+
// CHECK-NEXT: acc.copyout accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copy>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
53+
54+
#pragma acc data copyin(parmVar)
55+
;
56+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
57+
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
58+
// CHECK-NEXT: acc.terminator
59+
// CHECK-NEXT: } loc
60+
// CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "parmVar"}
61+
#pragma acc data copyin(parmVar) copyin(localVar1)
62+
;
63+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
64+
// CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"}
65+
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
66+
// CHECK-NEXT: acc.terminator
67+
// CHECK-NEXT: } loc
68+
// CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "localVar1"}
69+
// CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "parmVar"}
70+
#pragma acc data copyin(parmVar, localVar1)
71+
;
72+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
73+
// CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"}
74+
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
75+
// CHECK-NEXT: acc.terminator
76+
// CHECK-NEXT: } loc
77+
// CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "localVar1"}
78+
// CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, name = "parmVar"}
79+
80+
#pragma acc data copyin(capture, always, alwaysin, readonly: parmVar)
81+
;
82+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier always,readonly,capture>, name = "parmVar"}
83+
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>) {
84+
// CHECK-NEXT: acc.terminator
85+
// CHECK-NEXT: } loc
86+
// CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always,readonly,capture>, name = "parmVar"}
87+
88+
#pragma acc data copyin(capture: parmVar) copyin(readonly, always, alwaysin: parmVar)
89+
;
90+
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
91+
// CHECK-NEXT: %[[COPYIN2:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier always,readonly>, name = "parmVar"}
92+
// CHECK-NEXT: acc.data dataOperands(%[[COPYIN1]], %[[COPYIN2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
93+
// CHECK-NEXT: acc.terminator
94+
// CHECK-NEXT: } loc
95+
// CHECK-NEXT: acc.delete accPtr(%[[COPYIN2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier always,readonly>, name = "parmVar"}
96+
// CHECK-NEXT: acc.delete accPtr(%[[COPYIN1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_copyin>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
97+
98+
#pragma acc data copyout(parmVar)
99+
;
100+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar"}
101+
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>) {
102+
// CHECK-NEXT: acc.terminator
103+
// CHECK-NEXT: } loc
104+
// CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar"}
105+
106+
#pragma acc data copyout(parmVar) copyout(localVar1)
107+
;
108+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar"}
109+
// CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "localVar1"}
110+
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
111+
// CHECK-NEXT: acc.terminator
112+
// CHECK-NEXT: } loc
113+
// CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) {name = "localVar1"}
114+
// CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar"}
115+
116+
#pragma acc data copyout(parmVar, localVar1)
117+
;
118+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar"}
119+
// CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "localVar1"}
120+
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
121+
// CHECK-NEXT: acc.terminator
122+
// CHECK-NEXT: } loc
123+
// CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[LV1]] : !cir.ptr<!s32i>) {name = "localVar1"}
124+
// CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar"}
125+
126+
#pragma acc data copyout(capture, zero, alwaysout: parmVar)
127+
;
128+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout,capture>, name = "parmVar"}
129+
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>) {
130+
// CHECK-NEXT: acc.terminator
131+
// CHECK-NEXT: } loc
132+
// CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout,capture>, name = "parmVar"}
133+
134+
#pragma acc data copyout(capture: parmVar) copyout(always, alwaysout: parmVar)
135+
;
136+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
137+
// CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier always>, name = "parmVar"}
138+
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
139+
// CHECK-NEXT: acc.terminator
140+
// CHECK-NEXT: } loc
141+
// CHECK-NEXT: acc.copyout accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier always>, name = "parmVar"}
142+
// CHECK-NEXT: acc.copyout accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
143+
144+
#pragma acc data create(parmVar)
145+
;
146+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
147+
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>) {
148+
// CHECK-NEXT: acc.terminator
149+
// CHECK-NEXT: } loc
150+
// CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "parmVar"}
151+
152+
#pragma acc data create(parmVar) create(localVar1)
153+
;
154+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
155+
// CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"}
156+
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
157+
// CHECK-NEXT: acc.terminator
158+
// CHECK-NEXT: } loc
159+
// CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "localVar1"}
160+
// CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "parmVar"}
161+
162+
#pragma acc data create(parmVar, localVar1)
163+
;
164+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar"}
165+
// CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[LV1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "localVar1"}
166+
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
167+
// CHECK-NEXT: acc.terminator
168+
// CHECK-NEXT: } loc
169+
// CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "localVar1"}
170+
// CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, name = "parmVar"}
171+
172+
173+
#pragma acc data create(capture, zero: parmVar)
174+
;
175+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero,capture>, name = "parmVar"}
176+
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>) {
177+
// CHECK-NEXT: acc.terminator
178+
// CHECK-NEXT: } loc
179+
// CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero,capture>, name = "parmVar"}
180+
181+
#pragma acc data create(capture: parmVar) create(zero: parmVar)
182+
;
183+
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
184+
// CHECK-NEXT: %[[CREATE2:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "parmVar"}
185+
// CHECK-NEXT: acc.data dataOperands(%[[CREATE1]], %[[CREATE2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
186+
// CHECK-NEXT: acc.terminator
187+
// CHECK-NEXT: } loc
188+
// CHECK-NEXT: acc.delete accPtr(%[[CREATE2]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier zero>, name = "parmVar"}
189+
// CHECK-NEXT: acc.delete accPtr(%[[CREATE1]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_create>, modifiers = #acc<data_clause_modifier capture>, name = "parmVar"}
190+
}

0 commit comments

Comments
 (0)