Skip to content

Commit 125dbe1

Browse files
erichkeaneerichkeane
authored
[OpenACC][CIR] 'update' construct lowering + a few clauses (#146378)
The 'update' construct has 3 'var-list' clauses, device, self, and host. Each has a pretty simple data-operand type syntax in the IR, so this patch implements them as well. At least one of those is required to be present on an 'update', so we cannot do any lowering without them. Note that 'self' and 'host' are aliases.
1 parent 2d825cc commit 125dbe1

File tree

3 files changed

+102
-5
lines changed

3 files changed

+102
-5
lines changed

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

Lines changed: 31 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -613,12 +613,39 @@ class OpenACCClauseCIREmitter final
613613
} else {
614614
llvm_unreachable("var-list version of self shouldn't get here");
615615
}
616+
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
617+
assert(!clause.isEmptySelfClause() && !clause.isConditionExprClause() &&
618+
"var-list version of self required for update");
619+
for (const Expr *var : clause.getVarList())
620+
addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
621+
var, mlir::acc::DataClause::acc_update_self, {},
622+
/*structured=*/false, /*implicit=*/false);
616623
} else if constexpr (isCombinedType<OpTy>) {
617624
applyToComputeOp(clause);
618625
} else {
619-
// TODO: When we've implemented this for everything, switch this to an
620-
// unreachable. update construct remains.
621-
return clauseNotImplemented(clause);
626+
llvm_unreachable("Unknown construct kind in VisitSelfClause");
627+
}
628+
}
629+
630+
void VisitHostClause(const OpenACCHostClause &clause) {
631+
if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
632+
for (const Expr *var : clause.getVarList())
633+
addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
634+
var, mlir::acc::DataClause::acc_update_host, {},
635+
/*structured=*/false, /*implicit=*/false);
636+
} else {
637+
llvm_unreachable("Unknown construct kind in VisitHostClause");
638+
}
639+
}
640+
641+
void VisitDeviceClause(const OpenACCDeviceClause &clause) {
642+
if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
643+
for (const Expr *var : clause.getVarList())
644+
addDataOperand<mlir::acc::UpdateDeviceOp>(
645+
var, mlir::acc::DataClause::acc_update_device, {},
646+
/*structured=*/false, /*implicit=*/false);
647+
} else {
648+
llvm_unreachable("Unknown construct kind in VisitDeviceClause");
622649
}
623650
}
624651

@@ -1095,6 +1122,7 @@ EXPL_SPEC(mlir::acc::WaitOp)
10951122
EXPL_SPEC(mlir::acc::HostDataOp)
10961123
EXPL_SPEC(mlir::acc::EnterDataOp)
10971124
EXPL_SPEC(mlir::acc::ExitDataOp)
1125+
EXPL_SPEC(mlir::acc::UpdateOp)
10981126
#undef EXPL_SPEC
10991127

11001128
template <typename ComputeOp, typename LoopOp>

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -266,8 +266,10 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
266266

267267
mlir::LogicalResult
268268
CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
269-
cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");
270-
return mlir::failure();
269+
mlir::Location start = getLoc(s.getSourceRange().getBegin());
270+
emitOpenACCOp<UpdateOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
271+
s.clauses());
272+
return mlir::success();
271273
}
272274
mlir::LogicalResult
273275
CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {

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

Lines changed: 67 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,67 @@
1+
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
2+
3+
void acc_update(int parmVar, int *ptrParmVar) {
4+
// CHECK: cir.func{{.*}} @acc_update(%[[ARG:.*]]: !s32i{{.*}}, %[[PTRARG:.*]]: !cir.ptr<!s32i>{{.*}}) {
5+
// CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar", init]
6+
// CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["ptrParmVar", init]
7+
// CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr<!s32i>
8+
// CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
9+
10+
#pragma acc update device(parmVar)
11+
// CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
12+
// CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]] : !cir.ptr<!s32i>)
13+
14+
#pragma acc update device(parmVar, ptrParmVar)
15+
// CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
16+
// CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
17+
// CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]], %[[UPD_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
18+
19+
#pragma acc update device(parmVar) device(ptrParmVar)
20+
// CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
21+
// CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
22+
// CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]], %[[UPD_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
23+
24+
#pragma acc update host(parmVar)
25+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_host>, name = "parmVar", structured = false}
26+
// CHECK-NEXT: acc.update dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
27+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
28+
29+
#pragma acc update host(parmVar, ptrParmVar)
30+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_host>, name = "parmVar", structured = false}
31+
// CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_host>, name = "ptrParmVar", structured = false}
32+
// CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
33+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {name = "ptrParmVar", structured = false}
34+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
35+
36+
#pragma acc update host(parmVar) host(ptrParmVar)
37+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_host>, name = "parmVar", structured = false}
38+
// CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_host>, name = "ptrParmVar", structured = false}
39+
// CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
40+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {name = "ptrParmVar", structured = false}
41+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
42+
43+
#pragma acc update self(parmVar)
44+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
45+
// CHECK-NEXT: acc.update dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
46+
// 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}
47+
48+
#pragma acc update self(parmVar, ptrParmVar)
49+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
50+
// CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
51+
// CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
52+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
53+
// 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}
54+
55+
#pragma acc update self(parmVar) self(ptrParmVar)
56+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
57+
// CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
58+
// CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
59+
// CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
60+
// 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}
61+
62+
#pragma acc update self(parmVar) device(ptrParmVar)
63+
// CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
64+
// CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
65+
// CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[UPD_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
66+
// 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+
}

0 commit comments

Comments
 (0)