Skip to content

Commit 857815f

Browse files
erichkeaneerichkeane
authored
[OpenACC][CIR] Implement 'rest' of update clause lowering (#146414)
This implements the async, wait, if, and if_present (as well as device_type, but that is a detail of async/wait) lowering. All of these are implemented the same way they are for the compute constructs, so this is a pretty mild amount of changes.
1 parent e44fbea commit 857815f

File tree

4 files changed

+186
-18
lines changed

4 files changed

+186
-18
lines changed

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

Lines changed: 17 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -376,7 +376,8 @@ class OpenACCClauseCIREmitter final
376376
// on all operation types.
377377
mlir::ArrayAttr getAsyncOnlyAttr() {
378378
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
379-
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
379+
mlir::acc::KernelsOp, mlir::acc::DataOp,
380+
mlir::acc::UpdateOp>) {
380381
return operation.getAsyncOnlyAttr();
381382
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
382383
mlir::acc::ExitDataOp>) {
@@ -401,7 +402,8 @@ class OpenACCClauseCIREmitter final
401402
// on all operation types.
402403
mlir::ArrayAttr getAsyncOperandsDeviceTypeAttr() {
403404
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
404-
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
405+
mlir::acc::KernelsOp, mlir::acc::DataOp,
406+
mlir::acc::UpdateOp>) {
405407
return operation.getAsyncOperandsDeviceTypeAttr();
406408
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
407409
mlir::acc::ExitDataOp>) {
@@ -427,7 +429,8 @@ class OpenACCClauseCIREmitter final
427429
// on all operation types.
428430
mlir::OperandRange getAsyncOperands() {
429431
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
430-
mlir::acc::KernelsOp, mlir::acc::DataOp>)
432+
mlir::acc::KernelsOp, mlir::acc::DataOp,
433+
mlir::acc::UpdateOp>)
431434
return operation.getAsyncOperands();
432435
else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
433436
mlir::acc::ExitDataOp>)
@@ -522,7 +525,8 @@ class OpenACCClauseCIREmitter final
522525
decodeDeviceType(clause.getArchitectures()[0].getIdentifierInfo()));
523526
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
524527
mlir::acc::SerialOp, mlir::acc::KernelsOp,
525-
mlir::acc::DataOp, mlir::acc::LoopOp>) {
528+
mlir::acc::DataOp, mlir::acc::LoopOp,
529+
mlir::acc::UpdateOp>) {
526530
// Nothing to do here, these constructs don't have any IR for these, as
527531
// they just modify the other clauses IR. So setting of
528532
// `lastDeviceTypeValues` (done above) is all we need.
@@ -531,7 +535,7 @@ class OpenACCClauseCIREmitter final
531535
// 'lastDeviceTypeValues' to set the value for the child visitor.
532536
} else {
533537
// TODO: When we've implemented this for everything, switch this to an
534-
// unreachable. update, data, routine constructs remain.
538+
// unreachable. routine construct remains.
535539
return clauseNotImplemented(clause);
536540
}
537541
}
@@ -566,7 +570,8 @@ class OpenACCClauseCIREmitter final
566570
hasAsyncClause = true;
567571
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
568572
mlir::acc::KernelsOp, mlir::acc::DataOp,
569-
mlir::acc::EnterDataOp, mlir::acc::ExitDataOp>) {
573+
mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
574+
mlir::acc::UpdateOp>) {
570575
if (!clause.hasIntExpr()) {
571576
operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
572577
} else {
@@ -655,27 +660,20 @@ class OpenACCClauseCIREmitter final
655660
mlir::acc::ShutdownOp, mlir::acc::SetOp,
656661
mlir::acc::DataOp, mlir::acc::WaitOp,
657662
mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
658-
mlir::acc::ExitDataOp>) {
663+
mlir::acc::ExitDataOp, mlir::acc::UpdateOp>) {
659664
operation.getIfCondMutable().append(
660665
createCondition(clause.getConditionExpr()));
661666
} else if constexpr (isCombinedType<OpTy>) {
662667
applyToComputeOp(clause);
663668
} else {
664-
// 'if' applies to most of the constructs, but hold off on lowering them
665-
// until we can write tests/know what we're doing with codegen to make
666-
// sure we get it right.
667-
// TODO: When we've implemented this for everything, switch this to an
668-
// unreachable. update construct remains.
669-
return clauseNotImplemented(clause);
669+
llvm_unreachable("Unknown construct kind in VisitIfClause");
670670
}
671671
}
672672

673673
void VisitIfPresentClause(const OpenACCIfPresentClause &clause) {
674-
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
674+
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp,
675+
mlir::acc::UpdateOp>) {
675676
operation.setIfPresent(true);
676-
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
677-
// Last unimplemented one here, so just put it in this way instead.
678-
return clauseNotImplemented(clause);
679677
} else {
680678
llvm_unreachable("unknown construct kind in VisitIfPresentClause");
681679
}
@@ -710,7 +708,8 @@ class OpenACCClauseCIREmitter final
710708
void VisitWaitClause(const OpenACCWaitClause &clause) {
711709
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
712710
mlir::acc::KernelsOp, mlir::acc::DataOp,
713-
mlir::acc::EnterDataOp, mlir::acc::ExitDataOp>) {
711+
mlir::acc::EnterDataOp, mlir::acc::ExitDataOp,
712+
mlir::acc::UpdateOp>) {
714713
if (!clause.hasExprs()) {
715714
operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
716715
} else {

‎clang/test/CIR/CodeGenOpenACC/update.c

Lines changed: 111 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -64,4 +64,115 @@ void acc_update(int parmVar, int *ptrParmVar) {
6464
// CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
6565
// CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[UPD_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
6666
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
67+
68+
#pragma acc update self(parmVar) if (parmVar == 1)
69+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
70+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
71+
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1>
72+
// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
73+
// CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
74+
// CHECK-NEXT: acc.update if(%[[CMP_CAST]]) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
75+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
76+
#pragma acc update self(parmVar) if (parmVar == 1) if_present
77+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
78+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
79+
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1>
80+
// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
81+
// CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
82+
// CHECK-NEXT: acc.update if(%[[CMP_CAST]]) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>) attributes {ifPresent}
83+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
84+
85+
#pragma acc update self(parmVar) wait
86+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
87+
// CHECK-NEXT: acc.update wait dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
88+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
89+
90+
#pragma acc update self(parmVar) wait device_type(nvidia)
91+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
92+
// CHECK-NEXT: acc.update wait dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
93+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
94+
95+
#pragma acc update self(parmVar) device_type(radeon) wait
96+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
97+
// CHECK-NEXT: acc.update wait([#acc.device_type<radeon>]) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
98+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
99+
100+
#pragma acc update self(parmVar) wait(parmVar)
101+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
102+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
103+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
104+
// CHECK-NEXT: acc.update wait({%[[PARM_CAST]] : si32}) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
105+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
106+
107+
#pragma acc update self(parmVar) wait(parmVar) device_type(nvidia)
108+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
109+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
110+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
111+
// CHECK-NEXT: acc.update wait({%[[PARM_CAST]] : si32}) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
112+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
113+
114+
#pragma acc update self(parmVar) device_type(radeon) wait(parmVar)
115+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
116+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
117+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
118+
// CHECK-NEXT: acc.update wait({%[[PARM_CAST]] : si32} [#acc.device_type<radeon>]) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
119+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
120+
121+
#pragma acc update self(parmVar) device_type(radeon) wait(parmVar, 1, 2)
122+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
123+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
124+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
125+
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1>
126+
// CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]
127+
// CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2>
128+
// CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_CONST]]
129+
// CHECK-NEXT: acc.update wait({%[[PARM_CAST]] : si32, %[[ONE_CAST]] : si32, %[[TWO_CAST]] : si32} [#acc.device_type<radeon>]) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
130+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
131+
132+
#pragma acc update self(parmVar) device_type(radeon) wait(devnum:parmVar: 1, 2)
133+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
134+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
135+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
136+
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1>
137+
// CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]
138+
// CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2>
139+
// CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_CONST]]
140+
// CHECK-NEXT: acc.update wait({devnum: %[[PARM_CAST]] : si32, %[[ONE_CAST]] : si32, %[[TWO_CAST]] : si32} [#acc.device_type<radeon>]) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
141+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
142+
143+
#pragma acc update self(parmVar) async
144+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
145+
// CHECK-NEXT: acc.update async dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
146+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
147+
148+
#pragma acc update self(parmVar) async device_type(nvidia)
149+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
150+
// CHECK-NEXT: acc.update async dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
151+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
152+
153+
#pragma acc update self(parmVar) device_type(radeon) async
154+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async([#acc.device_type<radeon>]) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
155+
// CHECK-NEXT: acc.update async([#acc.device_type<radeon>]) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
156+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) async([#acc.device_type<radeon>]) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
157+
158+
#pragma acc update self(parmVar) async(parmVar)
159+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
160+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
161+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
162+
// CHECK-NEXT: acc.update async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
163+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
164+
165+
#pragma acc update self(parmVar) async(parmVar) device_type(nvidia)
166+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
167+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
168+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
169+
// CHECK-NEXT: acc.update async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
170+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
171+
172+
#pragma acc update self(parmVar) device_type(radeon) async(parmVar)
173+
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
174+
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
175+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32 [#acc.device_type<radeon>]) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
176+
// CHECK-NEXT: acc.update async(%[[PARM_CAST]] : si32 [#acc.device_type<radeon>]) dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
177+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32 [#acc.device_type<radeon>]) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
67178
}

‎mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3028,6 +3028,21 @@ def OpenACC_UpdateOp : OpenACC_Op<"update",
30283028
/// Return the wait devnum value clause for the given device_type if
30293029
/// present.
30303030
mlir::Value getWaitDevnum(mlir::acc::DeviceType deviceType);
3031+
/// Add an entry to the 'async-only' attribute (clause spelled without
3032+
/// arguments)for each of the additional device types (or a none if it is
3033+
/// empty).
3034+
void addAsyncOnly(MLIRContext *, llvm::ArrayRef<DeviceType>);
3035+
/// Add a value to the 'async' with the current list of device types.
3036+
void addAsyncOperand(MLIRContext *, mlir::Value,
3037+
llvm::ArrayRef<DeviceType>);
3038+
/// Add an entry to the 'wait-only' attribute (clause spelled without
3039+
/// arguments)for each of the additional device types (or a none if it is
3040+
/// empty).
3041+
void addWaitOnly(MLIRContext *, llvm::ArrayRef<DeviceType>);
3042+
/// Add an array-like entry to the 'wait' with the current list of device
3043+
/// types.
3044+
void addWaitOperands(MLIRContext *, bool hasDevnum, mlir::ValueRange,
3045+
llvm::ArrayRef<DeviceType>);
30313046
}];
30323047

30333048
let assemblyFormat = [{

0 commit comments

Comments
 (0)