Skip to content
Merged
Show file tree
Hide file tree
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
68 changes: 52 additions & 16 deletions clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -376,10 +376,19 @@ class OpenACCClauseCIREmitter final
// on all operation types.
mlir::ArrayAttr getAsyncOnlyAttr() {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp>)
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
return operation.getAsyncOnlyAttr();
else if constexpr (isCombinedType<OpTy>)
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
if (!operation.getAsyncAttr())
return mlir::ArrayAttr{};

llvm::SmallVector<mlir::Attribute> devTysTemp;
devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
builder.getContext(), mlir::acc::DeviceType::None));
return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
} else if constexpr (isCombinedType<OpTy>) {
return operation.computeOp.getAsyncOnlyAttr();
}

// Note: 'wait' has async as well, but it cannot have data clauses, so we
// don't have to handle them here.
Expand All @@ -391,10 +400,19 @@ class OpenACCClauseCIREmitter final
// on all operation types.
mlir::ArrayAttr getAsyncOperandsDeviceTypeAttr() {
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp>)
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
return operation.getAsyncOperandsDeviceTypeAttr();
else if constexpr (isCombinedType<OpTy>)
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
if (!operation.getAsyncOperand())
return mlir::ArrayAttr{};

llvm::SmallVector<mlir::Attribute> devTysTemp;
devTysTemp.push_back(mlir::acc::DeviceTypeAttr::get(
builder.getContext(), mlir::acc::DeviceType::None));
return mlir::ArrayAttr::get(builder.getContext(), devTysTemp);
} else if constexpr (isCombinedType<OpTy>) {
return operation.computeOp.getAsyncOperandsDeviceTypeAttr();
}

// Note: 'wait' has async as well, but it cannot have data clauses, so we
// don't have to handle them here.
Expand All @@ -409,6 +427,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>)
return operation.getAsyncOperandMutable();
else if constexpr (isCombinedType<OpTy>)
return operation.computeOp.getAsyncOperands();

Expand Down Expand Up @@ -542,10 +562,11 @@ class OpenACCClauseCIREmitter final
void VisitAsyncClause(const OpenACCAsyncClause &clause) {
hasAsyncClause = true;
if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
mlir::acc::KernelsOp, mlir::acc::DataOp>) {
if (!clause.hasIntExpr())
mlir::acc::KernelsOp, mlir::acc::DataOp,
mlir::acc::EnterDataOp>) {
if (!clause.hasIntExpr()) {
operation.addAsyncOnly(builder.getContext(), lastDeviceTypeValues);
else {
} else {

mlir::Value intExpr;
{
Expand All @@ -572,8 +593,8 @@ class OpenACCClauseCIREmitter final
applyToComputeOp(clause);
} else {
// TODO: When we've implemented this for everything, switch this to an
// unreachable. Combined constructs remain. Data, enter data, exit data,
// update constructs remain.
// unreachable. Combined constructs remain. Exit data,update constructs
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

nit: missing space

// remain.
return clauseNotImplemented(clause);
}
}
Expand Down Expand Up @@ -604,7 +625,7 @@ 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::HostDataOp, mlir::acc::EnterDataOp>) {
operation.getIfCondMutable().append(
createCondition(clause.getConditionExpr()));
} else if constexpr (isCombinedType<OpTy>) {
Expand Down Expand Up @@ -659,7 +680,8 @@ 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::KernelsOp, mlir::acc::DataOp,
mlir::acc::EnterDataOp>) {
if (!clause.hasExprs()) {
operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues);
} else {
Expand Down Expand Up @@ -866,11 +888,16 @@ class OpenACCClauseCIREmitter final
var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(),
/*structured=*/true,
/*implicit=*/false);
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::CopyinOp>(
var, mlir::acc::DataClause::acc_copyin, 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. enter-data, declare constructs remain.
// unreachable. declare construct remains.
return clauseNotImplemented(clause);
}
}
Expand Down Expand Up @@ -900,11 +927,16 @@ class OpenACCClauseCIREmitter final
var, mlir::acc::DataClause::acc_create, clause.getModifierList(),
/*structured=*/true,
/*implicit=*/false);
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::CreateOp>(
var, mlir::acc::DataClause::acc_create, 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. enter-data, declare constructs remain.
// unreachable. declare construct remains.
return clauseNotImplemented(clause);
}
}
Expand Down Expand Up @@ -974,12 +1006,15 @@ class OpenACCClauseCIREmitter final
addDataOperand<mlir::acc::AttachOp, mlir::acc::DetachOp>(
var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true,
/*implicit=*/false);
} else if constexpr (isOneOfTypes<OpTy, mlir::acc::EnterDataOp>) {
for (const Expr *var : clause.getVarList())
addDataOperand<mlir::acc::AttachOp>(
var, mlir::acc::DataClause::acc_attach, {},
/*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. enter data remains.
return clauseNotImplemented(clause);
llvm_unreachable("Unknown construct kind in VisitAttachClause");
}
}
};
Expand Down Expand Up @@ -1018,6 +1053,7 @@ EXPL_SPEC(mlir::acc::ShutdownOp)
EXPL_SPEC(mlir::acc::SetOp)
EXPL_SPEC(mlir::acc::WaitOp)
EXPL_SPEC(mlir::acc::HostDataOp)
EXPL_SPEC(mlir::acc::EnterDataOp)
#undef EXPL_SPEC

template <typename ComputeOp, typename LoopOp>
Expand Down
6 changes: 4 additions & 2 deletions clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -250,8 +250,10 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCHostDataConstruct(

mlir::LogicalResult CIRGenFunction::emitOpenACCEnterDataConstruct(
const OpenACCEnterDataConstruct &s) {
cgm.errorNYI(s.getSourceRange(), "OpenACC EnterData Construct");
return mlir::failure();
mlir::Location start = getLoc(s.getSourceRange().getBegin());
emitOpenACCOp<EnterDataOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
s.clauses());
return mlir::success();
}
mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
const OpenACCExitDataConstruct &s) {
Expand Down
125 changes: 125 additions & 0 deletions clang/test/CIR/CodeGenOpenACC/enter-data.c
Original file line number Diff line number Diff line change
@@ -0,0 +1,125 @@
// 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 enter data copyin(parmVar)
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)

#pragma acc enter data copyin(readonly, alwaysin: parmVar)
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)

#pragma acc enter data copyin(readonly, alwaysin: parmVar) async
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)

#pragma acc enter data async copyin(readonly, alwaysin: parmVar)
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)

#pragma acc enter data copyin(readonly, alwaysin: parmVar) async(parmVar)
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)

#pragma acc enter data async(parmVar) copyin(readonly, alwaysin: parmVar)
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
// CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier readonly,alwaysin>, name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)

#pragma acc enter data create(parmVar)
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)

#pragma acc enter data create(zero: parmVar)
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)

#pragma acc enter data create(zero: parmVar) async
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data async dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)

#pragma acc enter data create(zero: parmVar) async(parmVar)
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {modifiers = #acc<data_clause_modifier zero>, name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)

#pragma acc enter data attach(ptrParmVar)
// CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
// CHECK-NEXT: acc.enter_data dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>)

#pragma acc enter data attach(ptrParmVar) async
// CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) async -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
// CHECK-NEXT: acc.enter_data async dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>)

#pragma acc enter data attach(ptrParmVar) async(parmVar)
// CHECK-NEXT: %[[PARM_LOAD:.*]] = cir.load{{.*}} %[[PARM]]
// CHECK-NEXT: %[[PARM_CAST:.*]] = builtin.unrealized_conversion_cast %[[PARM_LOAD]]
// CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
// CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[ATTACH1]] : !cir.ptr<!cir.ptr<!s32i>>)

#pragma acc enter data if (parmVar == 1) copyin(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: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)

#pragma acc enter data async if (parmVar == 1) copyin(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: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) async dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)

#pragma acc enter data if (parmVar == 1) async(parmVar) copyin(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: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr<!s32i>) async(%[[PARM_CAST]] : si32) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr<!s32i>)

#pragma acc enter data wait create(parmVar)
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data wait dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)

#pragma acc enter data wait(1) create(parmVar)
// CHECK-NEXT: %[[ONE_CONST:.*]] = cir.const #cir.int<1> : !s32i
// CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE_CONST]]
// CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data wait(%[[ONE_CAST]] : si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)

#pragma acc enter data wait(parmVar, 1, 2) create(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: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data wait(%[[PARM_CAST]], %[[ONE_CAST]], %[[TWO_CAST]] : si32, si32, si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)

#pragma acc enter data wait(devnum: parmVar: 1, 2) create(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: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
// CHECK-NEXT: acc.enter_data wait_devnum(%[[PARM_CAST]] : si32) wait(%[[ONE_CAST]], %[[TWO_CAST]] : si32, si32) dataOperands(%[[CREATE1]] : !cir.ptr<!s32i>)

}
19 changes: 19 additions & 0 deletions mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
Original file line number Diff line number Diff line change
Expand Up @@ -2010,6 +2010,25 @@ def OpenACC_EnterDataOp : OpenACC_Op<"enter_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 = [{
Expand Down
Loading
Loading