https://github.com/erichkeane created
https://github.com/llvm/llvm-project/pull/146167
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.
>From e382471b54d4971c1668f39017af28e32198d7dc Mon Sep 17 00:00:00 2001
From: erichkeane <eke...@nvidia.com>
Date: Fri, 27 Jun 2025 14:53:15 -0700
Subject: [PATCH] [OpenACC][CIR] Implement 'exit data' construct + clauses
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.
---
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp | 65 +++++++--
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp | 8 +-
clang/test/CIR/CodeGenOpenACC/exit-data.c | 134 ++++++++++++++++++
.../mlir/Dialect/OpenACC/OpenACCOps.td | 20 +++
mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp | 47 ++++++
5 files changed, 260 insertions(+), 14 deletions(-)
create mode 100644 clang/test/CIR/CodeGenOpenACC/exit-data.c
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-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {name =
"parmVar", structured = false}
+
+#pragma acc exit data wait(devnum: parmVar: 1, 2) delete(parmVar)
+ // 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_devnum(%[[PARM_CAST]] : si32)
wait(%[[ONE_CAST]], %[[TWO_CAST]] : si32, si32) dataOperands(%[[GDP]] :
!cir.ptr<!s32i>)
+ // CHECK-NEXT: acc.delete accPtr(%[[GDP]] : !cir.ptr<!s32i>) {name =
"parmVar", structured = false}
+}
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index 3403e158c9f58..9aaf9040c25b7 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -2083,6 +2083,26 @@ def OpenACC_ExitDataOp : OpenACC_Op<"exit_data",
/// The i-th data operand passed.
Value getDataOperand(unsigned i);
+
+ /// Add an entry to the 'async-only' attribute (clause spelled without
+ /// arguments). DeviceType array is supplied even though it should always
be
+ /// empty, so this can mirror other versions of this function.
+ void addAsyncOnly(MLIRContext *, llvm::ArrayRef<DeviceType>);
+ /// Add a value to the 'async'. DeviceType array is supplied even though it
+ /// should always be empty, so this can mirror other versions of this
+ /// function.
+ void addAsyncOperand(MLIRContext *, mlir::Value,
+ llvm::ArrayRef<DeviceType>);
+
+ /// Add an entry to the 'wait-only' attribute (clause spelled without
+ /// arguments). DeviceType array is supplied even though it should always
be
+ /// empty, so this can mirror other versions of this function.
+ void addWaitOnly(MLIRContext *, llvm::ArrayRef<DeviceType>);
+ /// Add an array-like entry to the 'wait'. DeviceType array is supplied
+ /// even though it should always be empty, so this can mirror other
versions
+ /// of this function.
+ void addWaitOperands(MLIRContext *, bool hasDevnum, mlir::ValueRange,
+ llvm::ArrayRef<DeviceType>);
}];
let assemblyFormat = [{
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index f0516ef0f0f62..3cd3d2c52a923 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -3169,6 +3169,53 @@ void
ExitDataOp::getCanonicalizationPatterns(RewritePatternSet &results,
results.add<RemoveConstantIfCondition<ExitDataOp>>(context);
}
+void ExitDataOp::addAsyncOnly(
+ MLIRContext *context, llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ assert(effectiveDeviceTypes.empty());
+ assert(!getAsyncAttr());
+ assert(!getAsyncOperand());
+
+ setAsyncAttr(mlir::UnitAttr::get(context));
+}
+
+void ExitDataOp::addAsyncOperand(
+ MLIRContext *context, mlir::Value newValue,
+ llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ assert(effectiveDeviceTypes.empty());
+ assert(!getAsyncAttr());
+ assert(!getAsyncOperand());
+
+ getAsyncOperandMutable().append(newValue);
+}
+
+void ExitDataOp::addWaitOnly(MLIRContext *context,
+ llvm::ArrayRef<DeviceType> effectiveDeviceTypes)
{
+ assert(effectiveDeviceTypes.empty());
+ assert(!getWaitAttr());
+ assert(getWaitOperands().empty());
+ assert(!getWaitDevnum());
+
+ setWaitAttr(mlir::UnitAttr::get(context));
+}
+
+void ExitDataOp::addWaitOperands(
+ MLIRContext *context, bool hasDevnum, mlir::ValueRange newValues,
+ llvm::ArrayRef<DeviceType> effectiveDeviceTypes) {
+ assert(effectiveDeviceTypes.empty());
+ assert(!getWaitAttr());
+ assert(getWaitOperands().empty());
+ assert(!getWaitDevnum());
+
+ // if hasDevnum, the first value is the devnum. The 'rest' go into the
+ // operands list.
+ if (hasDevnum) {
+ getWaitDevnumMutable().append(newValues.front());
+ newValues = newValues.drop_front();
+ }
+
+ getWaitOperandsMutable().append(newValues);
+}
+
//===----------------------------------------------------------------------===//
// EnterDataOp
//===----------------------------------------------------------------------===//
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
