[clang] [OpenACC][CIR] Implement 'host_data' lowering, plus all clauses (PR #143136)
Erich Keane via cfe-commits
cfe-commits at lists.llvm.org
Fri Jun 6 06:17:55 PDT 2025
https://github.com/erichkeane created https://github.com/llvm/llvm-project/pull/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.
>From 8419c28ce1a8eb2c8a974dd38a011e23dd86403d Mon Sep 17 00:00:00 2001
From: erichkeane <ekeane at nvidia.com>
Date: Thu, 5 Jun 2025 15:12:00 -0700
Subject: [PATCH] [OpenACC][CIR] Implement 'host_data' lowering, plus all
clauses
'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.
---
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 41 +++++++++++++-
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 16 ++++--
clang/test/CIR/CodeGenOpenACC/host_data.c | 55 +++++++++++++++++++
3 files changed, 106 insertions(+), 6 deletions(-)
create mode 100644 clang/test/CIR/CodeGenOpenACC/host_data.c
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index f41f776225152..e3657e9014121 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -319,6 +319,21 @@ class OpenACCClauseCIREmitter final
dataOperands.push_back(afterOp.getOperation());
}
+ template <typename BeforeOpTy>
+ void addDataOperand(const Expr *varOperand, mlir::acc::DataClause dataClause,
+ bool structured, bool implicit) {
+ DataOperandInfo opInfo = getDataOperandInfo(dirKind, varOperand);
+ auto beforeOp =
+ builder.create<BeforeOpTy>(opInfo.beginLoc, opInfo.varValue, structured,
+ implicit, opInfo.name, opInfo.bounds);
+ operation.getDataClauseOperandsMutable().append(beforeOp.getResult());
+
+ // Set the 'rest' of the info for the operation.
+ beforeOp.setDataClause(dataClause);
+ // Make sure we record these, so 'async' values can be updated later.
+ dataOperands.push_back(beforeOp.getOperation());
+ }
+
// Helper function that covers for the fact that we don't have this function
// on all operation types.
mlir::ArrayAttr getAsyncOnlyAttr() {
@@ -550,7 +565,8 @@ class OpenACCClauseCIREmitter final
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::InitOp,
mlir::acc::ShutdownOp, mlir::acc::SetOp,
- mlir::acc::DataOp, mlir::acc::WaitOp>) {
+ mlir::acc::DataOp, mlir::acc::WaitOp,
+ mlir::acc::HostDataOp>) {
operation.getIfCondMutable().append(
createCondition(clause.getConditionExpr()));
} else if constexpr (isCombinedType<OpTy>) {
@@ -566,6 +582,17 @@ class OpenACCClauseCIREmitter final
}
}
+ void VisitIfPresentClause(const OpenACCIfPresentClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
+ operation.setIfPresent(true);
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
+ // Last unimplemented one here, so just put it in this way instead.
+ return clauseNotImplemented(clause);
+ } else {
+ llvm_unreachable("unknown construct kind in VisitIfPresentClause");
+ }
+ }
+
void VisitDeviceNumClause(const OpenACCDeviceNumClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::InitOp, mlir::acc::ShutdownOp,
mlir::acc::SetOp>) {
@@ -791,6 +818,17 @@ class OpenACCClauseCIREmitter final
return clauseNotImplemented(clause);
}
}
+
+ void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
+ for (auto var : clause.getVarList())
+ addDataOperand<mlir::acc::UseDeviceOp>(
+ var, mlir::acc::DataClause::acc_use_device,
+ /*structured=*/true, /*implicit=*/false);
+ } else {
+ llvm_unreachable("Unknown construct kind in VisitUseDeviceClause");
+ }
+ }
};
template <typename OpTy>
@@ -826,6 +864,7 @@ EXPL_SPEC(mlir::acc::InitOp)
EXPL_SPEC(mlir::acc::ShutdownOp)
EXPL_SPEC(mlir::acc::SetOp)
EXPL_SPEC(mlir::acc::WaitOp)
+EXPL_SPEC(mlir::acc::HostDataOp)
#undef EXPL_SPEC
template <typename ComputeOp, typename LoopOp>
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index d922ca0c74d5d..2aab9cecf93d8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -235,6 +235,17 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCCombinedConstruct(
llvm_unreachable("invalid compute construct kind");
}
}
+
+mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
+ const OpenACCHostDataConstruct &s) {
+ mlir::Location start = getLoc(s.getSourceRange().getBegin());
+ mlir::Location end = getLoc(s.getSourceRange().getEnd());
+
+ return emitOpenACCOpAssociatedStmt<HostDataOp, mlir::acc::TerminatorOp>(
+ start, end, s.getDirectiveKind(), s.getDirectiveLoc(), s.clauses(),
+ s.getStructuredBlock());
+}
+
mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
const OpenACCEnterDataConstruct &s) {
cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
@@ -245,11 +256,6 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
cgm.errorNYI(s.getSourceRange(), "OpenACC ExitData Construct");
return mlir::failure();
}
-mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(
- const OpenACCHostDataConstruct &s) {
- cgm.errorNYI(s.getSourceRange(), "OpenACC HostData Construct");
- return mlir::failure();
-}
mlir::LogicalResult
CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");
diff --git a/clang/test/CIR/CodeGenOpenACC/host_data.c b/clang/test/CIR/CodeGenOpenACC/host_data.c
new file mode 100644
index 0000000000000..4c3f7dd092a2f
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/host_data.c
@@ -0,0 +1,55 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_host_data(int cond, int var1, int var2) {
+ // CHECK: cir.func @acc_host_data(%[[ARG_COND:.*]]: !s32i {{.*}}, %[[ARG_V1:.*]]: !s32i {{.*}}, %[[ARG_V2:.*]]: !s32i {{.*}}) {
+ // CHECK-NEXT: %[[COND:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["cond", init]
+ // CHECK-NEXT: %[[V1:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["var1", init]
+ // CHECK-NEXT: %[[V2:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["var2", init]
+ // CHECK-NEXT: cir.store %[[ARG_COND]], %[[COND]] : !s32i, !cir.ptr<!s32i>
+ // CHECK-NEXT: cir.store %[[ARG_V1]], %[[V1]] : !s32i, !cir.ptr<!s32i>
+ // CHECK-NEXT: cir.store %[[ARG_V2]], %[[V2]] : !s32i, !cir.ptr<!s32i>
+
+#pragma acc host_data use_device(var1)
+ {}
+ // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+ // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]] : !cir.ptr<!s32i>) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+#pragma acc host_data use_device(var1, var2)
+ {}
+ // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+ // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+ // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+
+#pragma acc host_data use_device(var1, var2) if_present
+ {}
+ // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+ // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+ // CHECK-NEXT: acc.host_data dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } attributes {ifPresent}
+
+#pragma acc host_data use_device(var1, var2) if(cond)
+ {}
+ // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+ // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+ // CHECK-NEXT: %[[LOAD_COND:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[COND_BOOL:.*]] = cir.cast(int_to_bool, %[[LOAD_COND]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_BOOL]] : !cir.bool to i1
+ // CHECK-NEXT: acc.host_data if(%[[COND_CAST]]) dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } loc
+
+#pragma acc host_data use_device(var1, var2) if(cond) if_present
+ {}
+ // CHECK-NEXT: %[[USE_DEV1:.*]] = acc.use_device varPtr(%[[V1]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var1"}
+ // CHECK-NEXT: %[[USE_DEV2:.*]] = acc.use_device varPtr(%[[V2]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "var2"}
+ // CHECK-NEXT: %[[LOAD_COND:.*]] = cir.load{{.*}} %[[COND]] : !cir.ptr<!s32i>, !s32i
+ // CHECK-NEXT: %[[COND_BOOL:.*]] = cir.cast(int_to_bool, %[[LOAD_COND]] : !s32i), !cir.bool
+ // CHECK-NEXT: %[[COND_CAST:.*]] = builtin.unrealized_conversion_cast %[[COND_BOOL]] : !cir.bool to i1
+ // CHECK-NEXT: acc.host_data if(%[[COND_CAST]]) dataOperands(%[[USE_DEV1]], %[[USE_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!s32i>) {
+ // CHECK-NEXT: acc.terminator
+ // CHECK-NEXT: } attributes {ifPresent}
+}
More information about the cfe-commits
mailing list