caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[OpenACC][CIR] Implement 'exit data' construct + clauses (#146167)

Browse Source 
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.
This commit is contained in:
Erich Keane
2025-06-30 06:19:43 -07:00
committed by GitHub
parent 51dfe28f87
commit a99fee6989
5 changed files with 260 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

65
clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
View File

@@ -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 @@ public:
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 @@ public:
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 @@ public:
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 @@ public:
// 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 @@ public:
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 @@ public:
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 @@ public:
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 @@ public:
}
}
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>

8
clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
View File

@@ -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");

134
clang/test/CIR/CodeGenOpenACC/exit-data.c Normal file
View File

@@ -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}
}

20
mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
View File

@@ -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 = [{

47
mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
View File

@@ -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
//===----------------------------------------------------------------------===//