diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp index b17994ee8771..d982cc92d9b4 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp @@ -376,10 +376,19 @@ class OpenACCClauseCIREmitter final // on all operation types. mlir::ArrayAttr getAsyncOnlyAttr() { if constexpr (isOneOfTypes) + mlir::acc::KernelsOp, mlir::acc::DataOp>) { return operation.getAsyncOnlyAttr(); - else if constexpr (isCombinedType) + } else if constexpr (isOneOfTypes) { + if (!operation.getAsyncAttr()) + return mlir::ArrayAttr{}; + + llvm::SmallVector 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) { 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. @@ -391,10 +400,19 @@ class OpenACCClauseCIREmitter final // on all operation types. mlir::ArrayAttr getAsyncOperandsDeviceTypeAttr() { if constexpr (isOneOfTypes) + mlir::acc::KernelsOp, mlir::acc::DataOp>) { return operation.getAsyncOperandsDeviceTypeAttr(); - else if constexpr (isCombinedType) + } else if constexpr (isOneOfTypes) { + if (!operation.getAsyncOperand()) + return mlir::ArrayAttr{}; + + llvm::SmallVector 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) { 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. @@ -409,6 +427,8 @@ class OpenACCClauseCIREmitter final if constexpr (isOneOfTypes) return operation.getAsyncOperands(); + else if constexpr (isOneOfTypes) + return operation.getAsyncOperandMutable(); else if constexpr (isCombinedType) return operation.computeOp.getAsyncOperands(); @@ -542,10 +562,11 @@ public: void VisitAsyncClause(const OpenACCAsyncClause &clause) { hasAsyncClause = true; if constexpr (isOneOfTypes) { - 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; { @@ -572,8 +593,8 @@ public: 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 + // remain. return clauseNotImplemented(clause); } } @@ -604,7 +625,7 @@ public: 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) { @@ -659,7 +680,8 @@ public: void VisitWaitClause(const OpenACCWaitClause &clause) { if constexpr (isOneOfTypes) { + mlir::acc::KernelsOp, mlir::acc::DataOp, + mlir::acc::EnterDataOp>) { if (!clause.hasExprs()) { operation.addWaitOnly(builder.getContext(), lastDeviceTypeValues); } else { @@ -866,11 +888,16 @@ public: var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(), /*structured=*/true, /*implicit=*/false); + } else if constexpr (isOneOfTypes) { + for (const Expr *var : clause.getVarList()) + addDataOperand( + var, mlir::acc::DataClause::acc_copyin, clause.getModifierList(), + /*structured=*/false, /*implicit=*/false); } else if constexpr (isCombinedType) { 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); } } @@ -900,11 +927,16 @@ public: var, mlir::acc::DataClause::acc_create, clause.getModifierList(), /*structured=*/true, /*implicit=*/false); + } else if constexpr (isOneOfTypes) { + for (const Expr *var : clause.getVarList()) + addDataOperand( + var, mlir::acc::DataClause::acc_create, clause.getModifierList(), + /*structured=*/false, /*implicit=*/false); } else if constexpr (isCombinedType) { 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); } } @@ -974,12 +1006,15 @@ public: addDataOperand( var, mlir::acc::DataClause::acc_attach, {}, /*structured=*/true, /*implicit=*/false); + } else if constexpr (isOneOfTypes) { + for (const Expr *var : clause.getVarList()) + addDataOperand( + var, mlir::acc::DataClause::acc_attach, {}, + /*structured=*/false, /*implicit=*/false); } else if constexpr (isCombinedType) { 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"); } } }; @@ -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 diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp index 1feefa55eb27..10a5601476f4 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp @@ -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(start, s.getDirectiveKind(), s.getDirectiveLoc(), + s.clauses()); + return mlir::success(); } mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct( const OpenACCExitDataConstruct &s) { diff --git a/clang/test/CIR/CodeGenOpenACC/enter-data.c b/clang/test/CIR/CodeGenOpenACC/enter-data.c new file mode 100644 index 000000000000..1785fba1a105 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/enter-data.c @@ -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{{.*}}) { + // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr, ["parmVar", init] + // CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr, !cir.ptr>, ["ptrParmVar", init] + // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr + // CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr, !cir.ptr> + +#pragma acc enter data copyin(parmVar) + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data dataOperands(%[[COPYIN1]] : !cir.ptr) + +#pragma acc enter data copyin(readonly, alwaysin: parmVar) + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {modifiers = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data dataOperands(%[[COPYIN1]] : !cir.ptr) + +#pragma acc enter data copyin(readonly, alwaysin: parmVar) async + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr) async -> !cir.ptr {modifiers = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data async dataOperands(%[[COPYIN1]] : !cir.ptr) + +#pragma acc enter data async copyin(readonly, alwaysin: parmVar) + // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[PARM]] : !cir.ptr) async -> !cir.ptr {modifiers = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data async dataOperands(%[[COPYIN1]] : !cir.ptr) + +#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) async(%[[PARM_CAST]] : si32) -> !cir.ptr {modifiers = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr) + +#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) async(%[[PARM_CAST]] : si32) -> !cir.ptr {modifiers = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr) + +#pragma acc enter data create(parmVar) + // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data dataOperands(%[[CREATE1]] : !cir.ptr) + +#pragma acc enter data create(zero: parmVar) + // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {modifiers = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data dataOperands(%[[CREATE1]] : !cir.ptr) + +#pragma acc enter data create(zero: parmVar) async + // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr) async -> !cir.ptr {modifiers = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data async dataOperands(%[[CREATE1]] : !cir.ptr) + +#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) async(%[[PARM_CAST]] : si32) -> !cir.ptr {modifiers = #acc, name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[CREATE1]] : !cir.ptr) + +#pragma acc enter data attach(ptrParmVar) + // CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr>) -> !cir.ptr> {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.enter_data dataOperands(%[[ATTACH1]] : !cir.ptr>) + +#pragma acc enter data attach(ptrParmVar) async + // CHECK-NEXT: %[[ATTACH1:.*]] = acc.attach varPtr(%[[PTRPARM]] : !cir.ptr>) async -> !cir.ptr> {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.enter_data async dataOperands(%[[ATTACH1]] : !cir.ptr>) + +#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>) async(%[[PARM_CAST]] : si32) -> !cir.ptr> {name = "ptrParmVar", structured = false} + // CHECK-NEXT: acc.enter_data async(%[[PARM_CAST]] : si32) dataOperands(%[[ATTACH1]] : !cir.ptr>) + +#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) -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) dataOperands(%[[COPYIN1]] : !cir.ptr) + +#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) async -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) async dataOperands(%[[COPYIN1]] : !cir.ptr) + +#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) async(%[[PARM_CAST]] : si32) -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data if(%[[CMP_CAST]]) async(%[[PARM_CAST]] : si32) dataOperands(%[[COPYIN1]] : !cir.ptr) + +#pragma acc enter data wait create(parmVar) + // CHECK-NEXT: %[[CREATE1:.*]] = acc.create varPtr(%[[PARM]] : !cir.ptr) -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data wait dataOperands(%[[CREATE1]] : !cir.ptr) + +#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) -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data wait(%[[ONE_CAST]] : si32) dataOperands(%[[CREATE1]] : !cir.ptr) + +#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) -> !cir.ptr {name = "parmVar", structured = false} + // CHECK-NEXT: acc.enter_data wait(%[[PARM_CAST]], %[[ONE_CAST]], %[[TWO_CAST]] : si32, si32, si32) dataOperands(%[[CREATE1]] : !cir.ptr) + +#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) -> !cir.ptr {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) + +} diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td index 8cbdf710cfa6..3403e158c9f5 100644 --- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td +++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td @@ -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); + /// 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); + /// 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); + /// 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); }]; let assemblyFormat = [{ diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp index 37acb6acbfa9..f0516ef0f0f6 100644 --- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp +++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp @@ -3218,6 +3218,53 @@ void EnterDataOp::getCanonicalizationPatterns(RewritePatternSet &results, results.add>(context); } +void EnterDataOp::addAsyncOnly( + MLIRContext *context, llvm::ArrayRef effectiveDeviceTypes) { + assert(effectiveDeviceTypes.empty()); + assert(!getAsyncAttr()); + assert(!getAsyncOperand()); + + setAsyncAttr(mlir::UnitAttr::get(context)); +} + +void EnterDataOp::addAsyncOperand( + MLIRContext *context, mlir::Value newValue, + llvm::ArrayRef effectiveDeviceTypes) { + assert(effectiveDeviceTypes.empty()); + assert(!getAsyncAttr()); + assert(!getAsyncOperand()); + + getAsyncOperandMutable().append(newValue); +} + +void EnterDataOp::addWaitOnly(MLIRContext *context, + llvm::ArrayRef effectiveDeviceTypes) { + assert(effectiveDeviceTypes.empty()); + assert(!getWaitAttr()); + assert(getWaitOperands().empty()); + assert(!getWaitDevnum()); + + setWaitAttr(mlir::UnitAttr::get(context)); +} + +void EnterDataOp::addWaitOperands( + MLIRContext *context, bool hasDevnum, mlir::ValueRange newValues, + llvm::ArrayRef 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); +} + //===----------------------------------------------------------------------===// // AtomicReadOp //===----------------------------------------------------------------------===//