Skip to content

Commit c02403e

Browse files
erichkeaneerichkeane
authored
[OpenACC][CIR] Implement 'host_data' lowering, plus all clauses (#143136)
'host_data' has its own Op kind, so this handles the lowering there, it looks exactly like the other ones we've done so far, so nothing novel here. host_data takes 3 clauses, 1 of which is required. 'use_device' is required, and results in an acc.use_device operation, which then feeds into the dataOperands list on acc.host_data. 'if_present' is a simple attribute on the operand. 'if' is a condition on the operand, identical to our other handling of 'if'. This patch handles all of these.
1 parent 28e2256 commit c02403e

File tree

3 files changed

+106
-6
lines changed

3 files changed

+106
-6
lines changed

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

Lines changed: 40 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -319,6 +319,21 @@ class OpenACCClauseCIREmitter final
319319
dataOperands.push_back(afterOp.getOperation());
320320
}
321321

322+
template <typename BeforeOpTy>
323+
void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
324+
bool structured, bool implicit) {
325+
DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand);
326+
auto beforeOp =
327+
builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
328+
implicit, opInfo.name, opInfo.bounds);
329+
operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
330+
331+
// Set the 'rest' of the info for the operation.
332+
beforeOp.setDataClause(dataClause);
333+
// Make sure we record these, so 'async' values can be updated later.
334+
dataOperands.push_back(beforeOp.getOperation());
335+
}
336+
322337
// Helper function that covers for the fact that we don't have this function
323338
// on all operation types.
324339
mlir::ArrayAttr getAsyncOnlyAttr() {
@@ -550,7 +565,8 @@ class OpenACCClauseCIREmitter final
550565
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
551566
mlir::acc::KernelsOp, mlir::acc::InitOp,
552567
mlir::acc::ShutdownOp, mlir::acc::SetOp,
553-
mlir::acc::DataOp, mlir::acc::WaitOp>) {
568+
mlir::acc::DataOp, mlir::acc::WaitOp,
569+
mlir::acc::HostDataOp>) {
554570
operation.getIfCondMutable().append(
555571
createCondition(clause.getConditionExpr()));
556572
} else if constexpr (isCombinedType<OpTy>) {
@@ -566,6 +582,17 @@ class OpenACCClauseCIREmitter final
566582
}
567583
}
568584

585+
void VisitIfPresentClause(const OpenACCIfPresentClause &clause) {
586+
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
587+
operation.setIfPresent(true);
588+
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
589+
// Last unimplemented one here, so just put it in this way instead.
590+
return clauseNotImplemented(clause);
591+
} else {
592+
llvm_unreachable("unknown construct kind in VisitIfPresentClause");
593+
}
594+
}
595+
569596
void VisitDeviceNumClause(const OpenACCDeviceNumClause &clause) {
570597
if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
571598
mlir::acc::SetOp>) {
@@ -791,6 +818,17 @@ class OpenACCClauseCIREmitter final
791818
return clauseNotImplemented(clause);
792819
}
793820
}
821+
822+
void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
823+
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
824+
for (auto var : clause.getVarList())
825+
addDataOperand<mlir::acc::UseDeviceOp>(
826+
var, mlir::acc::DataClause::acc_use_device,
827+
/*structured=*/true, /*implicit=*/false);
828+
} else {
829+
llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
830+
}
831+
}
794832
};
795833

796834
template <typename OpTy>
@@ -826,6 +864,7 @@ EXPL_SPEC(mlir::acc::InitOp)
826864
EXPL_SPEC(mlir::acc::ShutdownOp)
827865
EXPL_SPEC(mlir::acc::SetOp)
828866
EXPL_SPEC(mlir::acc::WaitOp)
867+
EXPL_SPEC(mlir::acc::HostDataOp)
829868
#undef EXPL_SPEC
830869

831870
template <typename ComputeOp, typename LoopOp>

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

Lines changed: 11 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -235,6 +235,17 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct(
235235
llvm_unreachable("invalid compute construct kind");
236236
}
237237
}
238+
239+
mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
240+
const OpenACCHostDataConstruct &s) {
241+
mlir::Location start = getLoc(s.getSourceRange().getBegin());
242+
mlir::Location end = getLoc(s.getSourceRange().getEnd());
243+
244+
return emitOpenACCOpAssociatedStmt<HostDataOp, mlir::acc::TerminatorOp>(
245+
start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
246+
s.getStructuredBlock());
247+
}
248+
238249
mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
239250
const OpenACCEnterDataConstruct &s) {
240251
cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
@@ -245,11 +256,6 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
245256
cgm.errorNYI(s.getSourceRange(), "OpenACC ExitData Construct");
246257
return mlir::failure();
247258
}
248-
mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
249-
const OpenACCHostDataConstruct &s) {
250-
cgm.errorNYI(s.getSourceRange(), "OpenACC HostData Construct");
251-
return mlir::failure();
252-
}
253259
mlir::LogicalResult
254260
CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
255261
cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");

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

Lines changed: 55 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,55 @@
1+
// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
2+
3+
void acc_host_data(int cond, int var1, int var2) {
4+
// CHECK: cir.func @acc_host_data(%[[ARG_COND:.*]]: !s32i {{.*}}, %[[ARG_V1:.*]]: !s32i {{.*}}, %[[ARG_V2:.*]]: !s32i {{.*}}) {
5+
// CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
6+
// CHECK-NEXT: %[[V1:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["var1", init]
7+
// CHECK-NEXT: %[[V2:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["var2", init]
8+
// CHECK-NEXT: cir.store %[[ARG_COND]], %[[COND]] : !s32i, !cir.ptr<!s32i>
9+
// CHECK-NEXT: cir.store %[[ARG_V1]], %[[V1]] : !s32i, !cir.ptr<!s32i>
10+
// CHECK-NEXT: cir.store %[[ARG_V2]], %[[V2]] : !s32i, !cir.ptr<!s32i>
11+
12+
#pragma acc host_data use_device(var1)
13+
{}
14+
// CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
15+
// CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]] : !cir.ptr<!s32i>) {
16+
// CHECK-NEXT: acc.terminator
17+
// CHECK-NEXT: } loc
18+
#pragma acc host_data use_device(var1, var2)
19+
{}
20+
// CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
21+
// CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
22+
// CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
23+
// CHECK-NEXT: acc.terminator
24+
// CHECK-NEXT: } loc
25+
26+
#pragma acc host_data use_device(var1, var2) if_present
27+
{}
28+
// CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
29+
// CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
30+
// CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
31+
// CHECK-NEXT: acc.terminator
32+
// CHECK-NEXT: } attributes {ifPresent}
33+
34+
#pragma acc host_data use_device(var1, var2) if(cond)
35+
{}
36+
// CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
37+
// CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
38+
// CHECK-NEXT: %[[LOAD_COND:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i
39+
// CHECK-NEXT: %[[COND_BOOL:.*]] = cir.cast(int_to_bool, %[[LOAD_COND]] : !s32i), !cir.bool
40+
// CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_BOOL]] : !cir.bool to i1
41+
// CHECK-NEXT: acc.host_data if(%[[COND_CAST]]) dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
42+
// CHECK-NEXT: acc.terminator
43+
// CHECK-NEXT: } loc
44+
45+
#pragma acc host_data use_device(var1, var2) if(cond) if_present
46+
{}
47+
// CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
48+
// CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
49+
// CHECK-NEXT: %[[LOAD_COND:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i
50+
// CHECK-NEXT: %[[COND_BOOL:.*]] = cir.cast(int_to_bool, %[[LOAD_COND]] : !s32i), !cir.bool
51+
// CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_BOOL]] : !cir.bool to i1
52+
// CHECK-NEXT: acc.host_data if(%[[COND_CAST]]) dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
53+
// CHECK-NEXT: acc.terminator
54+
// CHECK-NEXT: } attributes {ifPresent}
55+
}

0 commit comments

Comments
 (0)