[clang] [mlir] [OpenACC][CIR] Implement 'exit data' construct + clauses (PR #146167)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Jun 27 15:03:24 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
@llvm/pr-subscribers-clangir
Author: Erich Keane (erichkeane)
<details>
<summary>Changes</summary>
Similar to 'enter data', except the data clauses have a 'getdeviceptr' operation before, so that they can properly use the 'exit' operation correctly. While this is a touch awkward, it fits perfectly into the existing infrastructure.
Same as with 'enter data', we had to add some add-functions for async and wait.
---
Patch is 24.62 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/146167.diff
5 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+53-12)
- (modified) clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp (+6-2)
- (added) clang/test/CIR/CodeGenOpenACC/exit-data.c (+134)
- (modified) mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td (+20)
- (modified) mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp (+47)
``````````diff
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index d982cc92d9b4b..cc0f3b77c1a65 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -378,7 +378,8 @@ class OpenACCClauseCIREmitter final
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
return operation.getAsyncOnlyAttr();
- } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
+ mlir::acc::ExitDataOp>) {
if (!operation.getAsyncAttr())
return mlir::ArrayAttr{};
@@ -402,7 +403,8 @@ class OpenACCClauseCIREmitter final
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
return operation.getAsyncOperandsDeviceTypeAttr();
- } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
+ mlir::acc::ExitDataOp>) {
if (!operation.getAsyncOperand())
return mlir::ArrayAttr{};
@@ -427,7 +429,8 @@ class OpenACCClauseCIREmitter final
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp>)
return operation.getAsyncOperands();
- else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>)
+ else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp,
+ mlir::acc::ExitDataOp>)
return operation.getAsyncOperandMutable();
else if constexpr (isCombinedType<OpTy>)
return operation.computeOp.getAsyncOperands();
@@ -563,7 +566,7 @@ class OpenACCClauseCIREmitter final
hasAsyncClause = true;
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp,
- mlir::acc::EnterDataOp>) {
+ mlir::acc::EnterDataOp, mlir::acc::ExitDataOp>) {
if (!clause.hasIntExpr()) {
operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
} else {
@@ -593,8 +596,7 @@ class OpenACCClauseCIREmitter final
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. Combined constructs remain. Exit data, update constructs
- // remain.
+ // unreachable. Combined constructs remain. update construct remains.
return clauseNotImplemented(clause);
}
}
@@ -625,7 +627,8 @@ class OpenACCClauseCIREmitter final
mlir::acc::KernelsOp, mlir::acc::InitOp,
mlir::acc::ShutdownOp, mlir::acc::SetOp,
mlir::acc::DataOp, mlir::acc::WaitOp,
- mlir::acc::HostDataOp, mlir::acc::EnterDataOp>) {
+ mlir::acc::HostDataOp, mlir::acc::EnterDataOp,
+ mlir::acc::ExitDataOp>) {
operation.getIfCondMutable().append(
createCondition(clause.getConditionExpr()));
} else if constexpr (isCombinedType<OpTy>) {
@@ -635,8 +638,7 @@ class OpenACCClauseCIREmitter final
// until we can write tests/know what we're doing with codegen to make
// sure we get it right.
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. Enter data, exit data, host_data, update constructs
- // remain.
+ // unreachable. update construct remains.
return clauseNotImplemented(clause);
}
}
@@ -681,7 +683,7 @@ class OpenACCClauseCIREmitter final
void VisitWaitClause(const OpenACCWaitClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp,
- mlir::acc::EnterDataOp>) {
+ mlir::acc::EnterDataOp, mlir::acc::ExitDataOp>) {
if (!clause.hasExprs()) {
operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
} else {
@@ -697,7 +699,7 @@ class OpenACCClauseCIREmitter final
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. Enter data, exit data, update constructs remain.
+ // unreachable. update construct remains.
return clauseNotImplemented(clause);
}
}
@@ -910,11 +912,17 @@ class OpenACCClauseCIREmitter final
var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
/*structured=*/true,
/*implicit=*/false);
+ } else if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::CopyoutOp>(
+ var, mlir::acc::DataClause::acc_copyout, clause.getModifierList(),
+ /*structured=*/false,
+ /*implicit=*/false);
} else if constexpr (isCombinedType<OpTy>) {
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
- // unreachable. exit data, declare constructs remain.
+ // unreachable. declare construct remains.
return clauseNotImplemented(clause);
}
}
@@ -941,6 +949,38 @@ class OpenACCClauseCIREmitter final
}
}
+ void VisitDeleteClause(const OpenACCDeleteClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DeleteOp>(
+ var, mlir::acc::DataClause::acc_delete, {},
+ /*structured=*/false,
+ /*implicit=*/false);
+ } else {
+ llvm_unreachable("Unknown construct kind in VisitDeleteClause");
+ }
+ }
+
+ void VisitDetachClause(const OpenACCDetachClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
+ for (const Expr *var : clause.getVarList())
+ addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::DetachOp>(
+ var, mlir::acc::DataClause::acc_detach, {},
+ /*structured=*/false,
+ /*implicit=*/false);
+ } else {
+ llvm_unreachable("Unknown construct kind in VisitDetachClause");
+ }
+ }
+
+ void VisitFinalizeClause(const OpenACCFinalizeClause &clause) {
+ if constexpr (isOneOfTypes<OpTy, mlir::acc::ExitDataOp>) {
+ operation.setFinalize(true);
+ } else {
+ llvm_unreachable("Unknown construct kind in VisitFinalizeClause");
+ }
+ }
+
void VisitUseDeviceClause(const OpenACCUseDeviceClause &clause) {
if constexpr (isOneOfTypes<OpTy, mlir::acc::HostDataOp>) {
for (const Expr *var : clause.getVarList())
@@ -1054,6 +1094,7 @@ EXPL_SPEC(mlir::acc::SetOp)
EXPL_SPEC(mlir::acc::WaitOp)
EXPL_SPEC(mlir::acc::HostDataOp)
EXPL_SPEC(mlir::acc::EnterDataOp)
+EXPL_SPEC(mlir::acc::ExitDataOp)
#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 10a5601476f4e..f3a635b7c83eb 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -255,11 +255,15 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
s.clauses());
return mlir::success();
}
+
mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
const OpenACCExitDataConstruct &s) {
- cgm.errorNYI(s.getSourceRange(), "OpenACC ExitData Construct");
- return mlir::failure();
+ mlir::Location start = getLoc(s.getSourceRange().getBegin());
+ emitOpenACCOp<ExitDataOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
+ s.clauses());
+ return mlir::success();
}
+
mlir::LogicalResult
CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");
diff --git a/clang/test/CIR/CodeGenOpenACC/exit-data.c b/clang/test/CIR/CodeGenOpenACC/exit-data.c
new file mode 100644
index 0000000000000..ff987d20d5b6c
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/exit-data.c
@@ -0,0 +1,134 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+void acc_data(int parmVar, int *ptrParmVar) {
+ // CHECK: cir.func{{.*}} @acc_data(%[[ARG:.*]]: !s32i{{.*}}, %[[PTRARG:.*]]: !cir.ptr<!s32i>{{.*}}) {
+ // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar", init]
+ // CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["ptrParmVar", init]
+ // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr<!s32i>
+ // CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
+
+#pragma acc exit data copyout(parmVar)
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data copyout(zero, alwaysout: parmVar)
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+
+#pragma acc exit data copyout(zero, alwaysout: parmVar) async
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+
+#pragma acc exit data async copyout(zero, alwaysout: parmVar)
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+
+#pragma acc exit data finalize copyout(zero, alwaysout: parmVar) async(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>) attributes {finalize}
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+
+#pragma acc exit data async(parmVar) copyout(zero, alwaysout: parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {modifiers = #acc<data_clause_modifier zero,alwaysout>, name = "parmVar", structured = false}
+
+#pragma acc exit data delete(parmVar) finalize
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_delete>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data dataOperands(%[[GDP]] : !cir.ptr<!s32i>) attributes {finalize}
+ // CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data delete(parmVar) async(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_delete>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) {name = "parmVar", structured = false}
+
+#pragma acc exit data detach(ptrParmVar)
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_detach>, name = "ptrParmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data dataOperands(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>)
+ // CHECK-NEXT: acc.detach accPtr(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>) {name = "ptrParmVar", structured = false}
+
+#pragma acc exit data detach(ptrParmVar) async
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_detach>, name = "ptrParmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async dataOperands(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>)
+ // CHECK-NEXT: acc.detach accPtr(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>) async {name = "ptrParmVar", structured = false}
+
+#pragma acc exit data detach(ptrParmVar) async(parmVar) finalize
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_detach>, name = "ptrParmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>) attributes {finalize}
+ // CHECK-NEXT: acc.detach accPtr(%[[GDP]] : !cir.ptr<!cir.ptr<!s32i>>) async(%[[PARM_CAST]] : si32) {name = "ptrParmVar", structured = false}
+
+#pragma acc exit data if (parmVar == 1) copyout(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
+ // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data if(%[[CMP_CAST]]) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data async if (parmVar == 1) copyout(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
+ // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data if(%[[CMP_CAST]]) async dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data if (parmVar == 1) async(parmVar) copyout(parmVar)
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[CMP:.*]] = cir.cmp(eq, %[[PARM_LOAD]], %[[ONE_CONST]])
+ // CHECK-NEXT: %[[CMP_CAST:.*]] = builtin.unrealized_conversion_cast %[[CMP]]
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copyout>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data if(%[[CMP_CAST]]) async(%[[PARM_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.copyout accPtr(%[[GDP]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data wait delete(parmVar)
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_delete>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data wait dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data wait(1) delete(parmVar)
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_delete>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data wait(%[[ONE_CAST]] : si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc exit data wait(parmVar, 1, 2) delete(parmVar) finalize
+ // CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
+ // CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
+ // CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
+ // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]
+ // CHECK-NEXT: %[[TWO_CONST:.*]] = cir.const #cir.int<2> : !s32i
+ // CHECK-NEXT: %[[TWO_CAST:.*]] = builtin.unrealized_conversion_cast %[[TWO_CONST]]
+ // CHECK-NEXT: %[[GDP:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_delete>, name = "parmVar", structured = false}
+ // CHECK-NEXT: acc.exit_data wait(%[[PARM_CAST]], %[[ONE_CAST]], %[[TWO_CAST]] : si32, si32, si32) dataOperands(%[[GDP]] : !cir.ptr<!s32i>) attributes {finalize}
+ // CHECK-N...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/146167
More information about the cfe-commits
mailing list