forked from OSchip/llvm-project
[mlir] Use value instead of getValue (NFC)
This commit is contained in:
parent
0e95921bc3
commit
c27d815249
|
@ -400,7 +400,7 @@ public:
|
|||
inline Value getValue(unsigned pos) const {
|
||||
assert(pos < getNumDimAndSymbolVars() && "Invalid position");
|
||||
assert(hasValue(pos) && "variable's Value not set");
|
||||
return values[pos].getValue();
|
||||
return values[pos].value();
|
||||
}
|
||||
|
||||
/// Returns true if the pos^th variable has an associated Value.
|
||||
|
|
|
@ -348,7 +348,7 @@ def LLVM_FNegOp : LLVM_UnaryFloatArithmeticOp<
|
|||
class MemoryOpWithAlignmentBase {
|
||||
code setAlignmentCode = [{
|
||||
if ($alignment.hasValue()) {
|
||||
auto align = $alignment.getValue();
|
||||
auto align = $alignment.value();
|
||||
if (align != 0)
|
||||
inst->setAlignment(llvm::Align(align));
|
||||
}
|
||||
|
|
|
@ -569,7 +569,7 @@ class MMA_SYNC_INTR {
|
|||
# " \"" # op[2].ptx_elt_type # "\" == eltypeC && "
|
||||
# " \"" # op[3].ptx_elt_type # "\" == eltypeD "
|
||||
# " && (sat.has_value() ? " # sat # " == static_cast<int>(*sat) : true)"
|
||||
# !if(!ne(b1op, ""), " && (b1Op.has_value() ? MMAB1Op::" # b1op # " == b1Op.getValue() : true)", "") # ")\n"
|
||||
# !if(!ne(b1op, ""), " && (b1Op.has_value() ? MMAB1Op::" # b1op # " == b1Op.value() : true)", "") # ")\n"
|
||||
# " return " #
|
||||
MMA_SYNC_NAME<layoutA, layoutB, b1op, sat, op[0], op[1], op[2], op[3]>.id # ";",
|
||||
"") // if supported
|
||||
|
@ -995,8 +995,8 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
|
|||
$shape.getM(), $shape.getN(), $shape.getK(),
|
||||
$b1Op, $intOverflowBehavior,
|
||||
$layoutA, $layoutB,
|
||||
$multiplicandAPtxType.getValue(),
|
||||
$multiplicandBPtxType.getValue(),
|
||||
$multiplicandAPtxType.value(),
|
||||
$multiplicandBPtxType.value(),
|
||||
op.accumPtxType(),
|
||||
op.resultPtxType());
|
||||
|
||||
|
|
|
@ -47,7 +47,7 @@ public:
|
|||
bool hasValue() const { return impl.has_value(); }
|
||||
|
||||
/// Access the internal ParseResult value.
|
||||
ParseResult getValue() const { return impl.getValue(); }
|
||||
ParseResult getValue() const { return impl.value(); }
|
||||
ParseResult operator*() const { return getValue(); }
|
||||
|
||||
private:
|
||||
|
|
|
@ -251,9 +251,9 @@ private:
|
|||
[callback = std::forward<FnT>(callback)](
|
||||
T type, SmallVectorImpl<Type> &results, ArrayRef<Type>) {
|
||||
if (Optional<Type> resultOpt = callback(type)) {
|
||||
bool wasSuccess = static_cast<bool>(resultOpt.getValue());
|
||||
bool wasSuccess = static_cast<bool>(resultOpt.value());
|
||||
if (wasSuccess)
|
||||
results.push_back(resultOpt.getValue());
|
||||
results.push_back(resultOpt.value());
|
||||
return Optional<LogicalResult>(success(wasSuccess));
|
||||
}
|
||||
return Optional<LogicalResult>();
|
||||
|
|
|
@ -2029,7 +2029,7 @@ IntegerRelation::unionBoundingBox(const IntegerRelation &otherCst) {
|
|||
if (!constLb.has_value() || !constOtherLb.has_value())
|
||||
return failure();
|
||||
std::fill(minLb.begin(), minLb.end(), 0);
|
||||
minLb.back() = std::min(constLb.getValue(), constOtherLb.getValue());
|
||||
minLb.back() = std::min(constLb.value(), constOtherLb.value());
|
||||
}
|
||||
|
||||
// Do the same for ub's but max of upper bounds. Identify max.
|
||||
|
@ -2045,7 +2045,7 @@ IntegerRelation::unionBoundingBox(const IntegerRelation &otherCst) {
|
|||
if (!constUb.has_value() || !constOtherUb.has_value())
|
||||
return failure();
|
||||
std::fill(maxUb.begin(), maxUb.end(), 0);
|
||||
maxUb.back() = std::max(constUb.getValue(), constOtherUb.getValue());
|
||||
maxUb.back() = std::max(constUb.value(), constOtherUb.value());
|
||||
}
|
||||
|
||||
std::fill(newLb.begin(), newLb.end(), 0);
|
||||
|
|
|
@ -646,7 +646,7 @@ public:
|
|||
IntegerAttr executionModeAttr = op.execution_modeAttr();
|
||||
std::string moduleName;
|
||||
if (module.getName().has_value())
|
||||
moduleName = "_" + module.getName().getValue().str();
|
||||
moduleName = "_" + module.getName().value().str();
|
||||
else
|
||||
moduleName = "";
|
||||
std::string executionModeInfoName =
|
||||
|
@ -1585,10 +1585,10 @@ void mlir::encodeBindAttribute(ModuleOp module) {
|
|||
if (descriptorSet && binding) {
|
||||
// Encode these numbers into the variable's symbolic name. If the
|
||||
// SPIR-V module has a name, add it at the beginning.
|
||||
auto moduleAndName = spvModule.getName().has_value()
|
||||
? spvModule.getName().getValue().str() + "_" +
|
||||
op.sym_name().str()
|
||||
: op.sym_name().str();
|
||||
auto moduleAndName =
|
||||
spvModule.getName().has_value()
|
||||
? spvModule.getName().value().str() + "_" + op.sym_name().str()
|
||||
: op.sym_name().str();
|
||||
std::string name =
|
||||
llvm::formatv("{0}_descriptor_set{1}_binding{2}", moduleAndName,
|
||||
std::to_string(descriptorSet.getInt()),
|
||||
|
|
|
@ -147,8 +147,8 @@ createLinalgBodyCalculationForElementwiseOp(Operation *op, ValueRange args,
|
|||
cast<tosa::NegateOp>(op).quantization_info()) {
|
||||
auto quantizationInfo = cast<tosa::NegateOp>(op).quantization_info();
|
||||
int32_t inputBitWidth = elementTy.getIntOrFloatBitWidth();
|
||||
int64_t inZp = quantizationInfo.getValue().getInputZp();
|
||||
int64_t outZp = quantizationInfo.getValue().getOutputZp();
|
||||
int64_t inZp = quantizationInfo.value().getInputZp();
|
||||
int64_t outZp = quantizationInfo.value().getOutputZp();
|
||||
|
||||
// Compute the maximum value that can occur in the intermediate buffer.
|
||||
int64_t zpAdd = inZp + outZp;
|
||||
|
@ -1166,7 +1166,7 @@ public:
|
|||
checkHasDynamicBatchDims(rewriter, op, {input, op.output()});
|
||||
if (!dynamicDimsOr.has_value())
|
||||
return failure();
|
||||
SmallVector<Value> dynamicDims = dynamicDimsOr.getValue();
|
||||
SmallVector<Value> dynamicDims = dynamicDimsOr.value();
|
||||
|
||||
// The shift and multiplier values.
|
||||
SmallVector<int32_t> multiplierValues;
|
||||
|
@ -1358,7 +1358,7 @@ public:
|
|||
checkHasDynamicBatchDims(rewriter, op, {input, op.output()});
|
||||
if (!dynamicDimsOr.has_value())
|
||||
return failure();
|
||||
SmallVector<Value> dynamicDims = dynamicDimsOr.getValue();
|
||||
SmallVector<Value> dynamicDims = dynamicDimsOr.value();
|
||||
|
||||
if (op.mode() != "NEAREST_NEIGHBOR" && op.mode() != "BILINEAR")
|
||||
return failure();
|
||||
|
@ -2053,7 +2053,7 @@ public:
|
|||
checkHasDynamicBatchDims(rewriter, op, {input, indices, op.output()});
|
||||
if (!dynamicDimsOr.has_value())
|
||||
return failure();
|
||||
SmallVector<Value> dynamicDims = dynamicDimsOr.getValue();
|
||||
SmallVector<Value> dynamicDims = dynamicDimsOr.value();
|
||||
|
||||
auto resultElementTy = resultTy.getElementType();
|
||||
|
||||
|
|
|
@ -696,7 +696,7 @@ public:
|
|||
checkHasDynamicBatchDims(rewriter, op, {input, op.output()});
|
||||
if (!dynamicDimsOr.has_value())
|
||||
return failure();
|
||||
SmallVector<Value> dynamicDims = dynamicDimsOr.getValue();
|
||||
SmallVector<Value> dynamicDims = dynamicDimsOr.value();
|
||||
|
||||
// Determine what the initial value needs to be for the max pool op.
|
||||
Attribute initialAttr;
|
||||
|
@ -773,7 +773,7 @@ public:
|
|||
checkHasDynamicBatchDims(rewriter, op, {input, op.output()});
|
||||
if (!dynamicDimsOr.has_value())
|
||||
return failure();
|
||||
SmallVector<Value> dynamicDims = dynamicDimsOr.getValue();
|
||||
SmallVector<Value> dynamicDims = dynamicDimsOr.value();
|
||||
|
||||
// Apply padding as necessary.
|
||||
llvm::SmallVector<int64_t> pad;
|
||||
|
|
|
@ -94,8 +94,8 @@ static void getXferIndices(OpBuilder &b, OpTy xferOp, Value iv,
|
|||
if (!isBroadcast) {
|
||||
AffineExpr d0, d1;
|
||||
bindDims(xferOp.getContext(), d0, d1);
|
||||
Value offset = adaptor.getIndices()[dim.getValue()];
|
||||
indices[dim.getValue()] =
|
||||
Value offset = adaptor.getIndices()[dim.value()];
|
||||
indices[dim.value()] =
|
||||
makeComposedAffineApply(b, loc, d0 + d1, {offset, iv});
|
||||
}
|
||||
}
|
||||
|
|
|
@ -1011,16 +1011,16 @@ void FlatAffineValueConstraints::getSliceBounds(
|
|||
auto ubConst = getConstantBound(BoundType::UB, pos);
|
||||
if (lbConst.has_value() && ubConst.has_value()) {
|
||||
// Detect equality to a constant.
|
||||
if (lbConst.getValue() == ubConst.getValue()) {
|
||||
memo[pos] = getAffineConstantExpr(lbConst.getValue(), context);
|
||||
if (lbConst.value() == ubConst.value()) {
|
||||
memo[pos] = getAffineConstantExpr(lbConst.value(), context);
|
||||
changed = true;
|
||||
continue;
|
||||
}
|
||||
|
||||
// Detect an variable as modulo of another variable w.r.t a
|
||||
// constant.
|
||||
if (detectAsMod(*this, pos, lbConst.getValue(), ubConst.getValue(),
|
||||
memo, context)) {
|
||||
if (detectAsMod(*this, pos, lbConst.value(), ubConst.value(), memo,
|
||||
context)) {
|
||||
changed = true;
|
||||
continue;
|
||||
}
|
||||
|
@ -1121,9 +1121,9 @@ void FlatAffineValueConstraints::getSliceBounds(
|
|||
<< "WARNING: Potentially over-approximating slice lb\n");
|
||||
auto lbConst = getConstantBound(BoundType::LB, pos + offset);
|
||||
if (lbConst.has_value()) {
|
||||
lbMap = AffineMap::get(
|
||||
numMapDims, numMapSymbols,
|
||||
getAffineConstantExpr(lbConst.getValue(), context));
|
||||
lbMap =
|
||||
AffineMap::get(numMapDims, numMapSymbols,
|
||||
getAffineConstantExpr(lbConst.value(), context));
|
||||
}
|
||||
}
|
||||
if (!ubMap || ubMap.getNumResults() > 1) {
|
||||
|
@ -1131,10 +1131,9 @@ void FlatAffineValueConstraints::getSliceBounds(
|
|||
<< "WARNING: Potentially over-approximating slice ub\n");
|
||||
auto ubConst = getConstantBound(BoundType::UB, pos + offset);
|
||||
if (ubConst.has_value()) {
|
||||
ubMap =
|
||||
AffineMap::get(numMapDims, numMapSymbols,
|
||||
getAffineConstantExpr(
|
||||
ubConst.getValue() + ubAdjustment, context));
|
||||
ubMap = AffineMap::get(
|
||||
numMapDims, numMapSymbols,
|
||||
getAffineConstantExpr(ubConst.value() + ubAdjustment, context));
|
||||
}
|
||||
}
|
||||
}
|
||||
|
@ -1674,12 +1673,12 @@ void FlatAffineRelation::compose(const FlatAffineRelation &other) {
|
|||
// Add and match domain of `rel` to domain of `this`.
|
||||
for (unsigned i = 0, e = rel.getNumDomainDims(); i < e; ++i)
|
||||
if (relMaybeValues[i].has_value())
|
||||
setValue(i, relMaybeValues[i].getValue());
|
||||
setValue(i, relMaybeValues[i].value());
|
||||
// Add and match range of `this` to range of `rel`.
|
||||
for (unsigned i = 0, e = getNumRangeDims(); i < e; ++i) {
|
||||
unsigned rangeIdx = rel.getNumDomainDims() + i;
|
||||
if (thisMaybeValues[rangeIdx].has_value())
|
||||
rel.setValue(rangeIdx, thisMaybeValues[rangeIdx].getValue());
|
||||
rel.setValue(rangeIdx, thisMaybeValues[rangeIdx].value());
|
||||
}
|
||||
|
||||
// Append `this` to `rel` and simplify constraints.
|
||||
|
|
|
@ -93,7 +93,7 @@ Optional<uint64_t> mlir::getConstantTripCount(AffineForOp forOp) {
|
|||
for (auto resultExpr : map.getResults()) {
|
||||
if (auto constExpr = resultExpr.dyn_cast<AffineConstantExpr>()) {
|
||||
if (tripCount.has_value())
|
||||
tripCount = std::min(tripCount.getValue(),
|
||||
tripCount = std::min(tripCount.value(),
|
||||
static_cast<uint64_t>(constExpr.getValue()));
|
||||
else
|
||||
tripCount = constExpr.getValue();
|
||||
|
@ -133,12 +133,12 @@ uint64_t mlir::getLargestDivisorOfTripCount(AffineForOp forOp) {
|
|||
thisGcd = resultExpr.getLargestKnownDivisor();
|
||||
}
|
||||
if (gcd.has_value())
|
||||
gcd = llvm::GreatestCommonDivisor64(gcd.getValue(), thisGcd);
|
||||
gcd = llvm::GreatestCommonDivisor64(gcd.value(), thisGcd);
|
||||
else
|
||||
gcd = thisGcd;
|
||||
}
|
||||
assert(gcd.has_value() && "value expected per above logic");
|
||||
return gcd.getValue();
|
||||
return gcd.value();
|
||||
}
|
||||
|
||||
/// Given an induction variable `iv` of type AffineForOp and an access `index`
|
||||
|
|
|
@ -376,7 +376,7 @@ Optional<int64_t> MemRefRegion::getConstantBoundingSizeAndShape(
|
|||
Optional<int64_t> diff =
|
||||
cstWithShapeBounds.getConstantBoundOnDimSize(d, &lb, &lbDivisor);
|
||||
if (diff.has_value()) {
|
||||
diffConstant = diff.getValue();
|
||||
diffConstant = diff.value();
|
||||
assert(diffConstant >= 0 && "Dim size bound can't be negative");
|
||||
assert(lbDivisor > 0);
|
||||
} else {
|
||||
|
@ -1013,7 +1013,7 @@ bool mlir::buildSliceTripCountMap(
|
|||
}
|
||||
Optional<uint64_t> maybeConstTripCount = getConstantTripCount(forOp);
|
||||
if (maybeConstTripCount.has_value()) {
|
||||
(*tripCountMap)[op] = maybeConstTripCount.getValue();
|
||||
(*tripCountMap)[op] = maybeConstTripCount.value();
|
||||
continue;
|
||||
}
|
||||
return false;
|
||||
|
@ -1022,7 +1022,7 @@ bool mlir::buildSliceTripCountMap(
|
|||
// Slice bounds are created with a constant ub - lb difference.
|
||||
if (!tripCount.has_value())
|
||||
return false;
|
||||
(*tripCountMap)[op] = tripCount.getValue();
|
||||
(*tripCountMap)[op] = tripCount.value();
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
@ -1322,7 +1322,7 @@ static Optional<int64_t> getMemoryFootprintBytes(Block &block,
|
|||
Optional<int64_t> size = region.second->getRegionSize();
|
||||
if (!size.has_value())
|
||||
return None;
|
||||
totalSizeInBytes += size.getValue();
|
||||
totalSizeInBytes += size.value();
|
||||
}
|
||||
return totalSizeInBytes;
|
||||
}
|
||||
|
|
|
@ -339,7 +339,7 @@ static bool isDimOpValidSymbol(OpTy dimOp, Region *region) {
|
|||
Optional<int64_t> index = dimOp.getConstantIndex();
|
||||
assert(index.has_value() &&
|
||||
"expect only `dim` operations with a constant index");
|
||||
int64_t i = index.getValue();
|
||||
int64_t i = index.value();
|
||||
return TypeSwitch<Operation *, bool>(dimOp.getSource().getDefiningOp())
|
||||
.Case<memref::ViewOp, memref::SubViewOp, memref::AllocOp>(
|
||||
[&](auto op) { return isMemRefSizeValidSymbol(op, i, region); })
|
||||
|
@ -1897,8 +1897,7 @@ struct AffineForEmptyLoopFolder : public OpRewritePattern<AffineForOp> {
|
|||
return failure();
|
||||
// Bail out when the loop iterates more than once and it returns any iterArg
|
||||
// out of order.
|
||||
if (tripCount.has_value() && tripCount.getValue() >= 2 &&
|
||||
iterArgsNotInOrder)
|
||||
if (tripCount.has_value() && tripCount.value() >= 2 && iterArgsNotInOrder)
|
||||
return failure();
|
||||
rewriter.replaceOp(forOp, replacements);
|
||||
return success();
|
||||
|
@ -1931,19 +1930,18 @@ OperandRange AffineForOp::getSuccessorEntryOperands(Optional<unsigned> index) {
|
|||
void AffineForOp::getSuccessorRegions(
|
||||
Optional<unsigned> index, ArrayRef<Attribute> operands,
|
||||
SmallVectorImpl<RegionSuccessor> ®ions) {
|
||||
assert((!index.has_value() || index.getValue() == 0) &&
|
||||
"expected loop region");
|
||||
assert((!index.has_value() || index.value() == 0) && "expected loop region");
|
||||
// The loop may typically branch back to its body or to the parent operation.
|
||||
// If the predecessor is the parent op and the trip count is known to be at
|
||||
// least one, branch into the body using the iterator arguments. And in cases
|
||||
// we know the trip count is zero, it can only branch back to its parent.
|
||||
Optional<uint64_t> tripCount = getTrivialConstantTripCount(*this);
|
||||
if (!index.has_value() && tripCount.has_value()) {
|
||||
if (tripCount.getValue() > 0) {
|
||||
if (tripCount.value() > 0) {
|
||||
regions.push_back(RegionSuccessor(&getLoopBody(), getRegionIterArgs()));
|
||||
return;
|
||||
}
|
||||
if (tripCount.getValue() == 0) {
|
||||
if (tripCount.value() == 0) {
|
||||
regions.push_back(RegionSuccessor(getResults()));
|
||||
return;
|
||||
}
|
||||
|
@ -3771,8 +3769,8 @@ ParseResult AffineParallelOp::parse(OpAsmParser &parser,
|
|||
arith::symbolizeAtomicRMWKind(attrVal.getValue());
|
||||
if (!reduction)
|
||||
return parser.emitError(loc, "invalid reduction value: ") << attrVal;
|
||||
reductions.push_back(builder.getI64IntegerAttr(
|
||||
static_cast<int64_t>(reduction.getValue())));
|
||||
reductions.push_back(
|
||||
builder.getI64IntegerAttr(static_cast<int64_t>(reduction.value())));
|
||||
// While we keep getting commas, keep parsing.
|
||||
return success();
|
||||
};
|
||||
|
|
|
@ -143,7 +143,7 @@ void AffineDataCopyGeneration::runOnBlock(Block *block,
|
|||
getMemoryFootprintBytes(forOp,
|
||||
/*memorySpace=*/0);
|
||||
return (footprint.has_value() &&
|
||||
static_cast<uint64_t>(footprint.getValue()) >
|
||||
static_cast<uint64_t>(footprint.value()) >
|
||||
fastMemCapacityBytes);
|
||||
};
|
||||
|
||||
|
|
|
@ -443,13 +443,13 @@ public:
|
|||
|
||||
if (firstSrcDepPos.has_value()) {
|
||||
if (lastDstDepPos.has_value()) {
|
||||
if (firstSrcDepPos.getValue() <= lastDstDepPos.getValue()) {
|
||||
if (firstSrcDepPos.value() <= lastDstDepPos.value()) {
|
||||
// No valid insertion point exists which preserves dependences.
|
||||
return nullptr;
|
||||
}
|
||||
}
|
||||
// Return the insertion point at 'firstSrcDepPos'.
|
||||
return depInsts[firstSrcDepPos.getValue()];
|
||||
return depInsts[firstSrcDepPos.value()];
|
||||
}
|
||||
// No dependence targets in range (or only dst deps in range), return
|
||||
// 'dstNodInst' insertion point.
|
||||
|
@ -942,10 +942,10 @@ static Value createPrivateMemRef(AffineForOp forOp, Operation *srcStoreOpInst,
|
|||
// Create 'newMemRefType' using 'newShape' from MemRefRegion accessed
|
||||
// by 'srcStoreOpInst'.
|
||||
uint64_t bufSize =
|
||||
getMemRefEltSizeInBytes(oldMemRefType) * numElements.getValue();
|
||||
getMemRefEltSizeInBytes(oldMemRefType) * numElements.value();
|
||||
unsigned newMemSpace;
|
||||
if (bufSize <= localBufSizeThreshold && fastMemorySpace.has_value()) {
|
||||
newMemSpace = fastMemorySpace.getValue();
|
||||
newMemSpace = fastMemorySpace.value();
|
||||
} else {
|
||||
newMemSpace = oldMemRefType.getMemorySpaceAsInt();
|
||||
}
|
||||
|
@ -1143,7 +1143,7 @@ static bool isFusionProfitable(Operation *srcOpInst, Operation *srcStoreOpInst,
|
|||
srcWriteRegion.getRegionSize();
|
||||
if (!maybeSrcWriteRegionSizeBytes.has_value())
|
||||
return false;
|
||||
int64_t srcWriteRegionSizeBytes = maybeSrcWriteRegionSizeBytes.getValue();
|
||||
int64_t srcWriteRegionSizeBytes = maybeSrcWriteRegionSizeBytes.value();
|
||||
|
||||
// Compute op instance count for the src loop nest.
|
||||
uint64_t dstLoopNestCost = getComputeCost(dstForOp, dstLoopNestStats);
|
||||
|
@ -1184,14 +1184,13 @@ static bool isFusionProfitable(Operation *srcOpInst, Operation *srcStoreOpInst,
|
|||
Optional<int64_t> maybeSliceWriteRegionSizeBytes =
|
||||
sliceWriteRegion.getRegionSize();
|
||||
if (!maybeSliceWriteRegionSizeBytes.has_value() ||
|
||||
maybeSliceWriteRegionSizeBytes.getValue() == 0) {
|
||||
maybeSliceWriteRegionSizeBytes.value() == 0) {
|
||||
LLVM_DEBUG(llvm::dbgs()
|
||||
<< "Failed to get slice write region size at loopDepth: " << i
|
||||
<< "\n");
|
||||
continue;
|
||||
}
|
||||
int64_t sliceWriteRegionSizeBytes =
|
||||
maybeSliceWriteRegionSizeBytes.getValue();
|
||||
int64_t sliceWriteRegionSizeBytes = maybeSliceWriteRegionSizeBytes.value();
|
||||
|
||||
// If we are fusing for reuse, check that write regions remain the same.
|
||||
// TODO: Write region check should check sizes and offsets in
|
||||
|
@ -1268,11 +1267,11 @@ static bool isFusionProfitable(Operation *srcOpInst, Operation *srcStoreOpInst,
|
|||
return false;
|
||||
}
|
||||
|
||||
auto srcMemSizeVal = srcMemSize.getValue();
|
||||
auto dstMemSizeVal = dstMemSize.getValue();
|
||||
auto srcMemSizeVal = srcMemSize.value();
|
||||
auto dstMemSizeVal = dstMemSize.value();
|
||||
|
||||
assert(sliceMemEstimate && "expected value");
|
||||
auto fusedMem = dstMemSizeVal + sliceMemEstimate.getValue();
|
||||
auto fusedMem = dstMemSizeVal + sliceMemEstimate.value();
|
||||
|
||||
LLVM_DEBUG(llvm::dbgs() << " src mem: " << srcMemSizeVal << "\n"
|
||||
<< " dst mem: " << dstMemSizeVal << "\n"
|
||||
|
|
|
@ -406,12 +406,12 @@ checkTilingLegalityImpl(MutableArrayRef<mlir::AffineForOp> origLoops) {
|
|||
for (unsigned k = 0, e = depComps.size(); k < e; k++) {
|
||||
DependenceComponent depComp = depComps[k];
|
||||
if (depComp.lb.has_value() && depComp.ub.has_value() &&
|
||||
depComp.lb.getValue() < depComp.ub.getValue() &&
|
||||
depComp.ub.getValue() < 0) {
|
||||
depComp.lb.value() < depComp.ub.value() &&
|
||||
depComp.ub.value() < 0) {
|
||||
LLVM_DEBUG(llvm::dbgs()
|
||||
<< "Dependence component lb = "
|
||||
<< Twine(depComp.lb.getValue())
|
||||
<< " ub = " << Twine(depComp.ub.getValue())
|
||||
<< Twine(depComp.lb.value())
|
||||
<< " ub = " << Twine(depComp.ub.value())
|
||||
<< " is negative at depth: " << Twine(d)
|
||||
<< " and thus violates the legality rule.\n");
|
||||
return false;
|
||||
|
@ -802,11 +802,11 @@ constructTiledIndexSetHyperRect(MutableArrayRef<AffineForOp> origLoops,
|
|||
newLoops[width + i].setStep(origLoops[i].getStep());
|
||||
|
||||
// Set the upper bound.
|
||||
if (mayBeConstantCount && mayBeConstantCount.getValue() < tileSizes[i]) {
|
||||
if (mayBeConstantCount && mayBeConstantCount.value() < tileSizes[i]) {
|
||||
// Trip count is less than the tile size: upper bound is lower bound +
|
||||
// trip count * stepSize.
|
||||
AffineMap ubMap = b.getSingleDimShiftAffineMap(
|
||||
mayBeConstantCount.getValue() * origLoops[i].getStep());
|
||||
mayBeConstantCount.value() * origLoops[i].getStep());
|
||||
newLoops[width + i].setUpperBound(
|
||||
/*operands=*/newLoops[i].getInductionVar(), ubMap);
|
||||
} else if (largestDiv % tileSizes[i] != 0) {
|
||||
|
@ -975,7 +975,7 @@ void mlir::getTileableBands(func::FuncOp f,
|
|||
LogicalResult mlir::loopUnrollFull(AffineForOp forOp) {
|
||||
Optional<uint64_t> mayBeConstantTripCount = getConstantTripCount(forOp);
|
||||
if (mayBeConstantTripCount.has_value()) {
|
||||
uint64_t tripCount = mayBeConstantTripCount.getValue();
|
||||
uint64_t tripCount = mayBeConstantTripCount.value();
|
||||
if (tripCount == 0)
|
||||
return success();
|
||||
if (tripCount == 1)
|
||||
|
@ -991,8 +991,8 @@ LogicalResult mlir::loopUnrollUpToFactor(AffineForOp forOp,
|
|||
uint64_t unrollFactor) {
|
||||
Optional<uint64_t> mayBeConstantTripCount = getConstantTripCount(forOp);
|
||||
if (mayBeConstantTripCount.has_value() &&
|
||||
mayBeConstantTripCount.getValue() < unrollFactor)
|
||||
return loopUnrollByFactor(forOp, mayBeConstantTripCount.getValue());
|
||||
mayBeConstantTripCount.value() < unrollFactor)
|
||||
return loopUnrollByFactor(forOp, mayBeConstantTripCount.value());
|
||||
return loopUnrollByFactor(forOp, unrollFactor);
|
||||
}
|
||||
|
||||
|
@ -1151,8 +1151,8 @@ LogicalResult mlir::loopUnrollJamUpToFactor(AffineForOp forOp,
|
|||
uint64_t unrollJamFactor) {
|
||||
Optional<uint64_t> mayBeConstantTripCount = getConstantTripCount(forOp);
|
||||
if (mayBeConstantTripCount.has_value() &&
|
||||
mayBeConstantTripCount.getValue() < unrollJamFactor)
|
||||
return loopUnrollJamByFactor(forOp, mayBeConstantTripCount.getValue());
|
||||
mayBeConstantTripCount.value() < unrollJamFactor)
|
||||
return loopUnrollJamByFactor(forOp, mayBeConstantTripCount.value());
|
||||
return loopUnrollJamByFactor(forOp, unrollJamFactor);
|
||||
}
|
||||
|
||||
|
@ -1574,7 +1574,7 @@ AffineForOp mlir::sinkSequentialLoops(AffineForOp forOp) {
|
|||
for (unsigned j = 0; j < maxLoopDepth; ++j) {
|
||||
DependenceComponent &depComp = depComps[j];
|
||||
assert(depComp.lb.has_value() && depComp.ub.has_value());
|
||||
if (depComp.lb.getValue() != 0 || depComp.ub.getValue() != 0)
|
||||
if (depComp.lb.value() != 0 || depComp.ub.value() != 0)
|
||||
isParallelLoop[j] = false;
|
||||
}
|
||||
}
|
||||
|
|
|
@ -1792,11 +1792,11 @@ MemRefType mlir::normalizeMemRefType(MemRefType memrefType, OpBuilder b,
|
|||
// For a static memref and an affine map with no symbols, this is
|
||||
// always bounded.
|
||||
assert(ubConst && "should always have an upper bound");
|
||||
if (ubConst.getValue() < 0)
|
||||
if (ubConst.value() < 0)
|
||||
// This is due to an invalid map that maps to a negative space.
|
||||
return memrefType;
|
||||
// If dimension of new memrefType is dynamic, the value is -1.
|
||||
newShape[d] = ubConst.getValue() + 1;
|
||||
newShape[d] = ubConst.value() + 1;
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -161,9 +161,9 @@ bufferization::getGlobalFor(arith::ConstantOp constantOp, uint64_t alignment) {
|
|||
if (!globalOp.getInitialValue().has_value())
|
||||
continue;
|
||||
uint64_t opAlignment = globalOp.getAlignment().has_value()
|
||||
? globalOp.getAlignment().getValue()
|
||||
? globalOp.getAlignment().value()
|
||||
: 0;
|
||||
Attribute initialValue = globalOp.getInitialValue().getValue();
|
||||
Attribute initialValue = globalOp.getInitialValue().value();
|
||||
if (opAlignment == alignment && initialValue == constantOp.getValue())
|
||||
return globalOp;
|
||||
}
|
||||
|
|
|
@ -116,7 +116,7 @@ static FuncOpAnalysisState getFuncOpAnalysisState(const AnalysisState &state,
|
|||
func::FuncDialect::getDialectNamespace());
|
||||
if (!maybeState.has_value())
|
||||
return FuncOpAnalysisState::NotAnalyzed;
|
||||
const auto &analyzedFuncOps = maybeState.getValue()->analyzedFuncOps;
|
||||
const auto &analyzedFuncOps = maybeState.value()->analyzedFuncOps;
|
||||
auto it = analyzedFuncOps.find(funcOp);
|
||||
if (it == analyzedFuncOps.end())
|
||||
return FuncOpAnalysisState::NotAnalyzed;
|
||||
|
|
|
@ -78,7 +78,7 @@ void gpu::SerializeToBlobPass::runOnOperation() {
|
|||
if (!maybeTargetISA.has_value())
|
||||
return signalPassFailure();
|
||||
|
||||
std::string targetISA = std::move(maybeTargetISA.getValue());
|
||||
std::string targetISA = std::move(maybeTargetISA.value());
|
||||
|
||||
LLVM_DEBUG({
|
||||
llvm::dbgs() << "ISA for module: " << getOperation().getNameAttr() << "\n";
|
||||
|
|
|
@ -239,7 +239,7 @@ ParseResult AllocaOp::parse(OpAsmParser &parser, OperationState &result) {
|
|||
result.attributes.getNamed("alignment");
|
||||
if (alignmentAttr.has_value()) {
|
||||
auto alignmentInt =
|
||||
alignmentAttr.getValue().getValue().dyn_cast<IntegerAttr>();
|
||||
alignmentAttr.value().getValue().dyn_cast<IntegerAttr>();
|
||||
if (!alignmentInt)
|
||||
return parser.emitError(parser.getNameLoc(),
|
||||
"expected integer alignment");
|
||||
|
@ -934,7 +934,7 @@ void InvokeOp::print(OpAsmPrinter &p) {
|
|||
|
||||
// Either function name or pointer
|
||||
if (isDirect)
|
||||
p.printSymbolName(callee.getValue());
|
||||
p.printSymbolName(callee.value());
|
||||
else
|
||||
p << getOperand(0);
|
||||
|
||||
|
@ -1239,7 +1239,7 @@ void CallOp::print(OpAsmPrinter &p) {
|
|||
// callee (first operand) otherwise.
|
||||
p << ' ';
|
||||
if (isDirect)
|
||||
p.printSymbolName(callee.getValue());
|
||||
p.printSymbolName(callee.value());
|
||||
else
|
||||
p << getOperand(0);
|
||||
|
||||
|
@ -2012,7 +2012,7 @@ LogicalResult GlobalOp::verify() {
|
|||
|
||||
Optional<uint64_t> alignAttr = getAlignment();
|
||||
if (alignAttr.has_value()) {
|
||||
uint64_t value = alignAttr.getValue();
|
||||
uint64_t value = alignAttr.value();
|
||||
if (!llvm::isPowerOf2_64(value))
|
||||
return emitError() << "alignment attribute is not a power of 2";
|
||||
}
|
||||
|
|
|
@ -118,14 +118,14 @@ MMATypes MmaOp::accumPtxType() {
|
|||
Optional<mlir::NVVM::MMATypes> val = inferOperandMMAType(
|
||||
getODSOperands(2).getTypes().front(), /*isAccum=*/true);
|
||||
assert(val.has_value() && "accumulator PTX type should always be inferrable");
|
||||
return val.getValue();
|
||||
return val.value();
|
||||
}
|
||||
|
||||
MMATypes MmaOp::resultPtxType() {
|
||||
Optional<mlir::NVVM::MMATypes> val =
|
||||
inferOperandMMAType(getResult().getType(), /*isAccum=*/true);
|
||||
assert(val.has_value() && "result PTX type should always be inferrable");
|
||||
return val.getValue();
|
||||
return val.value();
|
||||
}
|
||||
|
||||
void MmaOp::print(OpAsmPrinter &p) {
|
||||
|
@ -399,10 +399,10 @@ LogicalResult MmaOp::verify() {
|
|||
break;
|
||||
default:
|
||||
return emitError("invalid shape or multiplicand type: " +
|
||||
stringifyEnum(getMultiplicandAPtxType().getValue()));
|
||||
stringifyEnum(getMultiplicandAPtxType().value()));
|
||||
}
|
||||
|
||||
if (isIntegerPtxType(getMultiplicandAPtxType().getValue())) {
|
||||
if (isIntegerPtxType(getMultiplicandAPtxType().value())) {
|
||||
expectedResult.push_back(s32x4StructTy);
|
||||
expectedC.emplace_back(4, i32Ty);
|
||||
multiplicandFragType = i32Ty;
|
||||
|
@ -440,16 +440,16 @@ LogicalResult MmaOp::verify() {
|
|||
context, SmallVector<Type>(2, f64Ty)));
|
||||
allowedShapes.push_back({8, 8, 4});
|
||||
}
|
||||
if (isIntegerPtxType(getMultiplicandAPtxType().getValue())) {
|
||||
if (isIntegerPtxType(getMultiplicandAPtxType().value())) {
|
||||
expectedA.push_back({i32Ty});
|
||||
expectedB.push_back({i32Ty});
|
||||
expectedC.push_back({i32Ty, i32Ty});
|
||||
expectedResult.push_back(s32x2StructTy);
|
||||
if (isInt4PtxType(getMultiplicandAPtxType().getValue()))
|
||||
if (isInt4PtxType(getMultiplicandAPtxType().value()))
|
||||
allowedShapes.push_back({8, 8, 32});
|
||||
if (isInt8PtxType(getMultiplicandAPtxType().getValue()))
|
||||
if (isInt8PtxType(getMultiplicandAPtxType().value()))
|
||||
allowedShapes.push_back({8, 8, 16});
|
||||
if (getMultiplicandAPtxType().getValue() == MMATypes::b1)
|
||||
if (getMultiplicandAPtxType().value() == MMATypes::b1)
|
||||
allowedShapes.push_back({8, 8, 128});
|
||||
}
|
||||
}
|
||||
|
|
|
@ -777,16 +777,14 @@ fuseOperations(OpBuilder &b, LinalgOp rootOp, TiledLinalgOp tiledLinalgOp,
|
|||
LinalgDependenceGraph::DependenceType::RAW)
|
||||
continue;
|
||||
|
||||
unsigned resultIndex =
|
||||
dependence.getDependentOpViewResultNum().getValue();
|
||||
unsigned resultIndex = dependence.getDependentOpViewResultNum().value();
|
||||
LinalgOp consumer = origOpToFusedOp.lookup(dependence.getIndexingOp());
|
||||
if (!consumer)
|
||||
continue;
|
||||
|
||||
Value replacementValue = fusedOp.getOperation()->getResult(resultIndex);
|
||||
consumer.getOperation()->setOperand(
|
||||
dependence.getIndexingOpViewOperandNum().getValue(),
|
||||
replacementValue);
|
||||
dependence.getIndexingOpViewOperandNum().value(), replacementValue);
|
||||
}
|
||||
|
||||
// At this point, all Linalg uses of the tensors produced by `origOp` have
|
||||
|
|
|
@ -50,7 +50,7 @@ static Value allocBuffer(ImplicitLocOpBuilder &b,
|
|||
|
||||
IntegerAttr alignmentAttr;
|
||||
if (alignment.has_value())
|
||||
alignmentAttr = b.getI64IntegerAttr(alignment.getValue());
|
||||
alignmentAttr = b.getI64IntegerAttr(alignment.value());
|
||||
|
||||
// Static buffer.
|
||||
if (auto cst = allocSize.getDefiningOp<arith::ConstantIndexOp>()) {
|
||||
|
@ -234,7 +234,7 @@ FailureOr<PromotionInfo> mlir::linalg::promoteSubviewAsNewBuffer(
|
|||
Value size =
|
||||
failed(upperBound)
|
||||
? rangeValue.size
|
||||
: b.create<arith::ConstantIndexOp>(loc, upperBound.getValue());
|
||||
: b.create<arith::ConstantIndexOp>(loc, upperBound.value());
|
||||
LLVM_DEBUG(llvm::dbgs() << "Extracted tightest: " << size << "\n");
|
||||
fullSizes.push_back(size);
|
||||
partialSizes.push_back(
|
||||
|
|
|
@ -103,8 +103,7 @@ void mlir::linalg::LinalgTransformationFilter::
|
|||
replaceLinalgTransformationFilter(PatternRewriter &rewriter,
|
||||
Operation *op) const {
|
||||
if (replacement.has_value())
|
||||
op->setAttr(LinalgTransforms::kLinalgTransformMarker,
|
||||
replacement.getValue());
|
||||
op->setAttr(LinalgTransforms::kLinalgTransformMarker, replacement.value());
|
||||
else
|
||||
op->removeAttr(
|
||||
rewriter.getStringAttr(LinalgTransforms::kLinalgTransformMarker));
|
||||
|
@ -441,10 +440,10 @@ LogicalResult mlir::linalg::LinalgBaseTileAndFusePattern::matchAndRewrite(
|
|||
if (failed(unfusedTiledOp))
|
||||
return failure();
|
||||
rewriter.replaceOp(tiledAndFusedOps->op,
|
||||
getTiledOpResult(unfusedTiledOp.getValue()));
|
||||
getTiledOpResult(unfusedTiledOp.value()));
|
||||
tiledAndFusedOps->op = unfusedTiledOp->op;
|
||||
}
|
||||
op->replaceAllUsesWith(getTiledAndFusedOpResult(tiledAndFusedOps.getValue()));
|
||||
op->replaceAllUsesWith(getTiledAndFusedOpResult(tiledAndFusedOps.value()));
|
||||
|
||||
filter.replaceLinalgTransformationFilter(rewriter,
|
||||
tiledAndFusedOps->op.getOperation());
|
||||
|
|
|
@ -1301,7 +1301,7 @@ LogicalResult GlobalOp::verify() {
|
|||
// Verify that the initial value, if present, is either a unit attribute or
|
||||
// an elements attribute.
|
||||
if (getInitialValue().has_value()) {
|
||||
Attribute initValue = getInitialValue().getValue();
|
||||
Attribute initValue = getInitialValue().value();
|
||||
if (!initValue.isa<UnitAttr>() && !initValue.isa<ElementsAttr>())
|
||||
return emitOpError("initial value should be a unit or elements "
|
||||
"attribute, but got ")
|
||||
|
@ -1333,7 +1333,7 @@ LogicalResult GlobalOp::verify() {
|
|||
ElementsAttr GlobalOp::getConstantInitValue() {
|
||||
auto initVal = getInitialValue();
|
||||
if (getConstant() && initVal.has_value())
|
||||
return initVal.getValue().cast<ElementsAttr>();
|
||||
return initVal.value().cast<ElementsAttr>();
|
||||
return {};
|
||||
}
|
||||
|
||||
|
@ -2487,14 +2487,14 @@ static bool isTrivialSubViewOp(SubViewOp subViewOp) {
|
|||
// Check offsets are zero.
|
||||
if (llvm::any_of(mixedOffsets, [](OpFoldResult ofr) {
|
||||
Optional<int64_t> intValue = getConstantIntValue(ofr);
|
||||
return !intValue || intValue.getValue() != 0;
|
||||
return !intValue || intValue.value() != 0;
|
||||
}))
|
||||
return false;
|
||||
|
||||
// Check strides are one.
|
||||
if (llvm::any_of(mixedStrides, [](OpFoldResult ofr) {
|
||||
Optional<int64_t> intValue = getConstantIntValue(ofr);
|
||||
return !intValue || intValue.getValue() != 1;
|
||||
return !intValue || intValue.value() != 1;
|
||||
}))
|
||||
return false;
|
||||
|
||||
|
|
|
@ -340,11 +340,11 @@ scf::TileConsumerAndFuseProducersUsingSCFForOp::returningMatchAndRewrite(
|
|||
// 2c. Generate the tiled implementation of the producer of the source
|
||||
rewriter.setInsertionPoint(candidateSliceOp);
|
||||
FailureOr<Value> fusedProducerValue =
|
||||
tensor::replaceExtractSliceWithTiledProducer(
|
||||
rewriter, candidateSliceOp, fusableProducer.getValue());
|
||||
tensor::replaceExtractSliceWithTiledProducer(rewriter, candidateSliceOp,
|
||||
fusableProducer.value());
|
||||
if (failed(fusedProducerValue))
|
||||
continue;
|
||||
rewriter.replaceOp(candidateSliceOp, fusedProducerValue.getValue());
|
||||
rewriter.replaceOp(candidateSliceOp, fusedProducerValue.value());
|
||||
|
||||
// 2d. The operands of the fused producer might themselved be slices of
|
||||
// values produced by operations that implement the `TilingInterface`.
|
||||
|
|
|
@ -216,11 +216,11 @@ OpFoldResult spirv::LogicalAndOp::fold(ArrayRef<Attribute> operands) {
|
|||
|
||||
if (Optional<bool> rhs = getScalarOrSplatBoolAttr(operands.back())) {
|
||||
// x && true = x
|
||||
if (rhs.getValue())
|
||||
if (rhs.value())
|
||||
return operand1();
|
||||
|
||||
// x && false = false
|
||||
if (!rhs.getValue())
|
||||
if (!rhs.value())
|
||||
return operands.back();
|
||||
}
|
||||
|
||||
|
@ -247,12 +247,12 @@ OpFoldResult spirv::LogicalOrOp::fold(ArrayRef<Attribute> operands) {
|
|||
assert(operands.size() == 2 && "spv.LogicalOr should take two operands");
|
||||
|
||||
if (auto rhs = getScalarOrSplatBoolAttr(operands.back())) {
|
||||
if (rhs.getValue())
|
||||
if (rhs.value())
|
||||
// x || true = true
|
||||
return operands.back();
|
||||
|
||||
// x || false = x
|
||||
if (!rhs.getValue())
|
||||
if (!rhs.value())
|
||||
return operand1();
|
||||
}
|
||||
|
||||
|
|
|
@ -496,8 +496,8 @@ template <typename ParseType, typename... Args> struct ParseCommaSeparatedList {
|
|||
auto remainingValues = ParseCommaSeparatedList<Args...>{}(dialect, parser);
|
||||
if (!remainingValues)
|
||||
return llvm::None;
|
||||
return std::tuple_cat(std::tuple<ParseType>(parseVal.getValue()),
|
||||
remainingValues.getValue());
|
||||
return std::tuple_cat(std::tuple<ParseType>(parseVal.value()),
|
||||
remainingValues.value());
|
||||
}
|
||||
};
|
||||
|
||||
|
@ -607,11 +607,11 @@ static ParseResult parseStructMemberDecorations(
|
|||
|
||||
memberDecorationInfo.emplace_back(
|
||||
static_cast<uint32_t>(memberTypes.size() - 1), 1,
|
||||
memberDecoration.getValue(), memberDecorationValue.getValue());
|
||||
memberDecoration.value(), memberDecorationValue.value());
|
||||
} else {
|
||||
memberDecorationInfo.emplace_back(
|
||||
static_cast<uint32_t>(memberTypes.size() - 1), 0,
|
||||
memberDecoration.getValue(), 0);
|
||||
memberDecoration.value(), 0);
|
||||
}
|
||||
return success();
|
||||
};
|
||||
|
|
|
@ -131,8 +131,8 @@ static LogicalResult lowerEntryPointABIAttr(spirv::FuncOp funcOp,
|
|||
return funcOp.emitRemark("lower entry point failure: could not select "
|
||||
"execution model based on 'spv.target_env'");
|
||||
|
||||
builder.create<spirv::EntryPointOp>(
|
||||
funcOp.getLoc(), executionModel.getValue(), funcOp, interfaceVars);
|
||||
builder.create<spirv::EntryPointOp>(funcOp.getLoc(), executionModel.value(),
|
||||
funcOp, interfaceVars);
|
||||
|
||||
// Specifies the spv.ExecutionModeOp.
|
||||
auto localSizeAttr = entryPointAttr.getLocalSize();
|
||||
|
|
|
@ -1249,9 +1249,9 @@ OpFoldResult GetExtentOp::fold(ArrayRef<Attribute> operands) {
|
|||
Optional<int64_t> dim = getConstantDim();
|
||||
if (!dim.has_value())
|
||||
return nullptr;
|
||||
if (dim.getValue() >= elements.getNumElements())
|
||||
if (dim.value() >= elements.getNumElements())
|
||||
return nullptr;
|
||||
return elements.getValues<Attribute>()[(uint64_t)dim.getValue()];
|
||||
return elements.getValues<Attribute>()[(uint64_t)dim.value()];
|
||||
}
|
||||
|
||||
void GetExtentOp::build(OpBuilder &builder, OperationState &result, Value shape,
|
||||
|
|
|
@ -1770,7 +1770,7 @@ public:
|
|||
Optional<unsigned> optExp = merger.buildTensorExpFromLinalg(op);
|
||||
if (!optExp.has_value())
|
||||
return failure();
|
||||
unsigned exp = optExp.getValue();
|
||||
unsigned exp = optExp.value();
|
||||
|
||||
// Rejects an inadmissable tensor expression.
|
||||
OpOperand *sparseOut = nullptr;
|
||||
|
|
|
@ -894,7 +894,7 @@ Optional<unsigned> Merger::buildTensorExp(linalg::GenericOp op, Value v) {
|
|||
if (def->getNumOperands() == 1) {
|
||||
auto x = buildTensorExp(op, def->getOperand(0));
|
||||
if (x.has_value()) {
|
||||
unsigned e = x.getValue();
|
||||
unsigned e = x.value();
|
||||
if (isa<math::AbsOp>(def))
|
||||
return addExp(kAbsF, e);
|
||||
if (isa<complex::AbsOp>(def))
|
||||
|
@ -967,8 +967,8 @@ Optional<unsigned> Merger::buildTensorExp(linalg::GenericOp op, Value v) {
|
|||
auto x = buildTensorExp(op, def->getOperand(0));
|
||||
auto y = buildTensorExp(op, def->getOperand(1));
|
||||
if (x.has_value() && y.has_value()) {
|
||||
unsigned e0 = x.getValue();
|
||||
unsigned e1 = y.getValue();
|
||||
unsigned e0 = x.value();
|
||||
unsigned e1 = y.value();
|
||||
if (isa<arith::MulFOp>(def))
|
||||
return addExp(kMulF, e0, e1);
|
||||
if (isa<complex::MulOp>(def))
|
||||
|
|
|
@ -39,5 +39,5 @@ FailureOr<Value> tensor::replaceExtractSliceWithTiledProducer(
|
|||
if (failed(tiledResult))
|
||||
return failure();
|
||||
|
||||
return tiledResult.getValue();
|
||||
return tiledResult.value();
|
||||
}
|
||||
|
|
|
@ -201,7 +201,7 @@ public:
|
|||
rewriter, loc, weightPaddingAttr.getType(), weightPaddingAttr);
|
||||
|
||||
if (op.quantization_info().has_value()) {
|
||||
auto quantInfo = op.quantization_info().getValue();
|
||||
auto quantInfo = op.quantization_info().value();
|
||||
weight = createOpAndInfer<tosa::PadOp>(
|
||||
rewriter, loc, UnrankedTensorType::get(weightETy), weight,
|
||||
weightPaddingVal, nullptr,
|
||||
|
@ -265,7 +265,7 @@ public:
|
|||
rewriter, loc, inputPaddingAttr.getType(), inputPaddingAttr);
|
||||
|
||||
if (op.quantization_info().has_value()) {
|
||||
auto quantInfo = op.quantization_info().getValue();
|
||||
auto quantInfo = op.quantization_info().value();
|
||||
input = createOpAndInfer<tosa::PadOp>(
|
||||
rewriter, loc, UnrankedTensorType::get(inputETy), input,
|
||||
inputPaddingVal, nullptr,
|
||||
|
|
|
@ -238,7 +238,7 @@ LogicalResult mlir::reshapeLikeShapesAreCompatible(
|
|||
return emitError("invalid to have a single dimension (" +
|
||||
Twine(map.index()) +
|
||||
") expanded into multiple dynamic dims (" +
|
||||
Twine(expandedDimStart + dynamicShape.getValue()) +
|
||||
Twine(expandedDimStart + dynamicShape.value()) +
|
||||
"," + Twine(expandedDimStart + dim.index()) + ")");
|
||||
}
|
||||
dynamicShape = dim.index();
|
||||
|
|
|
@ -2783,8 +2783,8 @@ void TransferReadOp::build(OpBuilder &builder, OperationState &result,
|
|||
ValueRange indices, AffineMap permutationMap,
|
||||
Optional<ArrayRef<bool>> inBounds) {
|
||||
auto permutationMapAttr = AffineMapAttr::get(permutationMap);
|
||||
auto inBoundsAttr = (inBounds && !inBounds.getValue().empty())
|
||||
? builder.getBoolArrayAttr(inBounds.getValue())
|
||||
auto inBoundsAttr = (inBounds && !inBounds.value().empty())
|
||||
? builder.getBoolArrayAttr(inBounds.value())
|
||||
: ArrayAttr();
|
||||
build(builder, result, vectorType, source, indices, permutationMapAttr,
|
||||
inBoundsAttr);
|
||||
|
@ -2798,8 +2798,8 @@ void TransferReadOp::build(OpBuilder &builder, OperationState &result,
|
|||
AffineMap permutationMap = getTransferMinorIdentityMap(
|
||||
source.getType().cast<ShapedType>(), vectorType);
|
||||
auto permutationMapAttr = AffineMapAttr::get(permutationMap);
|
||||
auto inBoundsAttr = (inBounds && !inBounds.getValue().empty())
|
||||
? builder.getBoolArrayAttr(inBounds.getValue())
|
||||
auto inBoundsAttr = (inBounds && !inBounds.value().empty())
|
||||
? builder.getBoolArrayAttr(inBounds.value())
|
||||
: ArrayAttr();
|
||||
build(builder, result, vectorType, source, indices, permutationMapAttr,
|
||||
padding,
|
||||
|
@ -3322,8 +3322,8 @@ void TransferWriteOp::build(OpBuilder &builder, OperationState &result,
|
|||
AffineMap permutationMap,
|
||||
Optional<ArrayRef<bool>> inBounds) {
|
||||
auto permutationMapAttr = AffineMapAttr::get(permutationMap);
|
||||
auto inBoundsAttr = (inBounds && !inBounds.getValue().empty())
|
||||
? builder.getBoolArrayAttr(inBounds.getValue())
|
||||
auto inBoundsAttr = (inBounds && !inBounds.value().empty())
|
||||
? builder.getBoolArrayAttr(inBounds.value())
|
||||
: ArrayAttr();
|
||||
build(builder, result, vector, dest, indices, permutationMapAttr,
|
||||
/*mask=*/Value(), inBoundsAttr);
|
||||
|
|
|
@ -105,10 +105,9 @@ struct TransferReadPermutationLowering
|
|||
|
||||
// Transpose in_bounds attribute.
|
||||
ArrayAttr newInBoundsAttr =
|
||||
op.getInBounds()
|
||||
? transposeInBoundsAttr(rewriter, op.getInBounds().getValue(),
|
||||
permutation)
|
||||
: ArrayAttr();
|
||||
op.getInBounds() ? transposeInBoundsAttr(
|
||||
rewriter, op.getInBounds().value(), permutation)
|
||||
: ArrayAttr();
|
||||
|
||||
// Generate new transfer_read operation.
|
||||
VectorType newReadType =
|
||||
|
@ -176,10 +175,9 @@ struct TransferWritePermutationLowering
|
|||
|
||||
// Transpose in_bounds attribute.
|
||||
ArrayAttr newInBoundsAttr =
|
||||
op.getInBounds()
|
||||
? transposeInBoundsAttr(rewriter, op.getInBounds().getValue(),
|
||||
permutation)
|
||||
: ArrayAttr();
|
||||
op.getInBounds() ? transposeInBoundsAttr(
|
||||
rewriter, op.getInBounds().value(), permutation)
|
||||
: ArrayAttr();
|
||||
|
||||
// Generate new transfer_write operation.
|
||||
Value newVec = rewriter.create<vector::TransposeOp>(
|
||||
|
|
|
@ -555,7 +555,7 @@ public:
|
|||
kind, rewriter, isInt);
|
||||
if (!mult.has_value())
|
||||
return failure();
|
||||
rewriter.replaceOp(op, mult.getValue());
|
||||
rewriter.replaceOp(op, mult.value());
|
||||
return success();
|
||||
}
|
||||
|
||||
|
@ -573,7 +573,7 @@ public:
|
|||
createContractArithOp(loc, a, op.getRhs(), r, kind, rewriter, isInt);
|
||||
if (!m.has_value())
|
||||
return failure();
|
||||
result = rewriter.create<vector::InsertOp>(loc, resType, m.getValue(),
|
||||
result = rewriter.create<vector::InsertOp>(loc, resType, m.value(),
|
||||
result, pos);
|
||||
}
|
||||
rewriter.replaceOp(op, result);
|
||||
|
@ -1940,8 +1940,8 @@ ContractionOpLowering::lowerReduction(vector::ContractionOp op,
|
|||
return rewriter.notifyMatchFailure(op, [&](Diagnostic &diag) {
|
||||
diag << "expected iterIndex=" << iterIndex << "to map to a RHS dimension";
|
||||
});
|
||||
int64_t lhsIndex = lookupLhs.getValue();
|
||||
int64_t rhsIndex = lookupRhs.getValue();
|
||||
int64_t lhsIndex = lookupLhs.value();
|
||||
int64_t rhsIndex = lookupRhs.value();
|
||||
int64_t dimSize = lhsType.getDimSize(lhsIndex);
|
||||
if (dimSize != rhsType.getDimSize(rhsIndex))
|
||||
return rewriter.notifyMatchFailure(op, [&](Diagnostic &diag) {
|
||||
|
|
|
@ -107,13 +107,13 @@ verifyTypesAlongAllEdges(Operation *op, Optional<unsigned> sourceNo,
|
|||
auto printEdgeName = [&](InFlightDiagnostic &diag) -> InFlightDiagnostic & {
|
||||
diag << "from ";
|
||||
if (sourceNo)
|
||||
diag << "Region #" << sourceNo.getValue();
|
||||
diag << "Region #" << sourceNo.value();
|
||||
else
|
||||
diag << "parent operands";
|
||||
|
||||
diag << " to ";
|
||||
if (succRegionNo)
|
||||
diag << "Region #" << succRegionNo.getValue();
|
||||
diag << "Region #" << succRegionNo.value();
|
||||
else
|
||||
diag << "parent results";
|
||||
return diag;
|
||||
|
|
|
@ -332,11 +332,11 @@ AffineExpr AffineParser::parseSymbolSSAIdExpr() {
|
|||
/// affine-expr ::= integer-literal
|
||||
AffineExpr AffineParser::parseIntegerExpr() {
|
||||
auto val = getToken().getUInt64IntegerValue();
|
||||
if (!val.has_value() || (int64_t)val.getValue() < 0)
|
||||
if (!val.has_value() || (int64_t)val.value() < 0)
|
||||
return emitError("constant too large for index"), nullptr;
|
||||
|
||||
consumeToken(Token::integer);
|
||||
return builder.getAffineConstantExpr((int64_t)val.getValue());
|
||||
return builder.getAffineConstantExpr((int64_t)val.value());
|
||||
}
|
||||
|
||||
/// Parses an expression that can be a valid operand of an affine expression.
|
||||
|
|
|
@ -680,7 +680,7 @@ DenseElementsAttr TensorLiteralParser::getStringAttr(SMLoc loc,
|
|||
ShapedType type,
|
||||
Type eltTy) {
|
||||
if (hexStorage.has_value()) {
|
||||
auto stringValue = hexStorage.getValue().getStringValue();
|
||||
auto stringValue = hexStorage.value().getStringValue();
|
||||
return DenseStringElementsAttr::get(type, {stringValue});
|
||||
}
|
||||
|
||||
|
|
|
@ -122,7 +122,7 @@ ParseResult Parser::parseNameOrFileLineColLocation(LocationAttr &loc) {
|
|||
return emitError("expected integer column number in FileLineColLoc");
|
||||
consumeToken(Token::integer);
|
||||
|
||||
loc = FileLineColLoc::get(ctx, str, line.getValue(), column.getValue());
|
||||
loc = FileLineColLoc::get(ctx, str, line.value(), column.value());
|
||||
return success();
|
||||
}
|
||||
|
||||
|
|
|
@ -2021,7 +2021,7 @@ ParseResult OperationParser::parseRegionBody(Region ®ion, SMLoc startLoc,
|
|||
<< "previously referenced here";
|
||||
}
|
||||
Location loc = entryArg.sourceLoc.has_value()
|
||||
? entryArg.sourceLoc.getValue()
|
||||
? entryArg.sourceLoc.value()
|
||||
: getEncodedSourceLocation(argInfo.location);
|
||||
BlockArgument arg = block->addArgument(entryArg.type, loc);
|
||||
|
||||
|
|
|
@ -311,7 +311,7 @@ Type Parser::parseNonFunctionType() {
|
|||
auto width = getToken().getIntTypeBitwidth();
|
||||
if (!width.has_value())
|
||||
return (emitError("invalid integer width"), nullptr);
|
||||
if (width.getValue() > IntegerType::kMaxWidth) {
|
||||
if (width.value() > IntegerType::kMaxWidth) {
|
||||
emitError(getToken().getLoc(), "integer bitwidth is limited to ")
|
||||
<< IntegerType::kMaxWidth << " bits";
|
||||
return nullptr;
|
||||
|
|
|
@ -239,7 +239,7 @@ StringRef AttrOrTypeParameter::getComparator() const {
|
|||
StringRef AttrOrTypeParameter::getCppType() const {
|
||||
if (auto *stringType = dyn_cast<llvm::StringInit>(getDef()))
|
||||
return stringType->getValue();
|
||||
return getDefValue<llvm::StringInit>("cppType").getValue();
|
||||
return getDefValue<llvm::StringInit>("cppType").value();
|
||||
}
|
||||
|
||||
StringRef AttrOrTypeParameter::getCppAccessorType() const {
|
||||
|
|
|
@ -433,7 +433,7 @@ GlobalOp Importer::processGlobal(llvm::GlobalVariable *gv) {
|
|||
uint64_t alignment = 0;
|
||||
llvm::MaybeAlign maybeAlign = gv->getAlign();
|
||||
if (maybeAlign.has_value()) {
|
||||
llvm::Align align = maybeAlign.getValue();
|
||||
llvm::Align align = maybeAlign.value();
|
||||
alignment = align.value();
|
||||
}
|
||||
|
||||
|
|
|
@ -30,7 +30,7 @@ static llvm::omp::ScheduleKind
|
|||
convertToScheduleKind(Optional<omp::ClauseScheduleKind> schedKind) {
|
||||
if (!schedKind.has_value())
|
||||
return llvm::omp::OMP_SCHEDULE_Default;
|
||||
switch (schedKind.getValue()) {
|
||||
switch (schedKind.value()) {
|
||||
case omp::ClauseScheduleKind::Static:
|
||||
return llvm::omp::OMP_SCHEDULE_Static;
|
||||
case omp::ClauseScheduleKind::Dynamic:
|
||||
|
|
|
@ -676,7 +676,7 @@ LogicalResult ModuleTranslation::convertGlobals() {
|
|||
|
||||
Optional<uint64_t> alignment = op.getAlignment();
|
||||
if (alignment.has_value())
|
||||
var->setAlignment(llvm::MaybeAlign(alignment.getValue()));
|
||||
var->setAlignment(llvm::MaybeAlign(alignment.value()));
|
||||
|
||||
globalsMapping.try_emplace(op, var);
|
||||
}
|
||||
|
|
|
@ -377,7 +377,7 @@ spirv::Deserializer::processFunction(ArrayRef<uint32_t> operands) {
|
|||
|
||||
std::string fnName = getFunctionSymbol(fnID);
|
||||
auto funcOp = opBuilder.create<spirv::FuncOp>(
|
||||
unknownLoc, fnName, functionType, fnControl.getValue());
|
||||
unknownLoc, fnName, functionType, fnControl.value());
|
||||
curFunction = funcMap[fnID] = funcOp;
|
||||
auto *entryBlock = funcOp.addEntryBlock();
|
||||
LLVM_DEBUG({
|
||||
|
@ -883,7 +883,7 @@ spirv::Deserializer::processCooperativeMatrixType(ArrayRef<uint32_t> operands) {
|
|||
unsigned columns = getConstantInt(operands[4]).getInt();
|
||||
|
||||
typeMap[operands[0]] = spirv::CooperativeMatrixNVType::get(
|
||||
elementTy, scope.getValue(), rows, columns);
|
||||
elementTy, scope.value(), rows, columns);
|
||||
return success();
|
||||
}
|
||||
|
||||
|
@ -1067,8 +1067,8 @@ spirv::Deserializer::processImageType(ArrayRef<uint32_t> operands) {
|
|||
<< operands[7];
|
||||
|
||||
typeMap[operands[0]] = spirv::ImageType::get(
|
||||
elementTy, dim.getValue(), depthInfo.getValue(), arrayedInfo.getValue(),
|
||||
samplingInfo.getValue(), samplerUseInfo.getValue(), format.getValue());
|
||||
elementTy, dim.value(), depthInfo.value(), arrayedInfo.value(),
|
||||
samplingInfo.value(), samplerUseInfo.value(), format.value());
|
||||
return success();
|
||||
}
|
||||
|
||||
|
|
|
@ -53,15 +53,13 @@ getDirectionVectorStr(bool ret, unsigned numCommonLoops, unsigned loopNestDepth,
|
|||
for (const auto &dependenceComponent : dependenceComponents) {
|
||||
std::string lbStr = "-inf";
|
||||
if (dependenceComponent.lb.has_value() &&
|
||||
dependenceComponent.lb.getValue() !=
|
||||
std::numeric_limits<int64_t>::min())
|
||||
lbStr = std::to_string(dependenceComponent.lb.getValue());
|
||||
dependenceComponent.lb.value() != std::numeric_limits<int64_t>::min())
|
||||
lbStr = std::to_string(dependenceComponent.lb.value());
|
||||
|
||||
std::string ubStr = "+inf";
|
||||
if (dependenceComponent.ub.has_value() &&
|
||||
dependenceComponent.ub.getValue() !=
|
||||
std::numeric_limits<int64_t>::max())
|
||||
ubStr = std::to_string(dependenceComponent.ub.getValue());
|
||||
dependenceComponent.ub.value() != std::numeric_limits<int64_t>::max())
|
||||
ubStr = std::to_string(dependenceComponent.ub.value());
|
||||
|
||||
result += "[" + lbStr + ", " + ubStr + "]";
|
||||
}
|
||||
|
|
|
@ -1418,7 +1418,7 @@ void RegionIfOp::getSuccessorRegions(
|
|||
SmallVectorImpl<RegionSuccessor> ®ions) {
|
||||
// We always branch to the join region.
|
||||
if (index.has_value()) {
|
||||
if (index.getValue() < 2)
|
||||
if (index.value() < 2)
|
||||
regions.push_back(RegionSuccessor(&getJoinRegion(), getJoinArgs()));
|
||||
else
|
||||
regions.push_back(RegionSuccessor(getResults()));
|
||||
|
|
|
@ -733,13 +733,13 @@ static LogicalResult generateNamedGenericOpOds(LinalgOpConfig &opConfig,
|
|||
assert(arg.indexAttrMap.has_value());
|
||||
assert(arg.defaultIndices.has_value());
|
||||
size_t size = arg.indexAttrMap->affineMap().getNumResults();
|
||||
assert(arg.defaultIndices.getValue().size() == size);
|
||||
assert(arg.defaultIndices.value().size() == size);
|
||||
static const char typeFmt[] = "RankedI64ElementsAttr<[{0}]>";
|
||||
static const char defFmt[] = "DefaultValuedAttr<{0}, \"{ {1} }\">:${2}";
|
||||
std::string defaultVals;
|
||||
llvm::raw_string_ostream ss(defaultVals);
|
||||
llvm::interleave(
|
||||
arg.defaultIndices.getValue(), ss,
|
||||
arg.defaultIndices.value(), ss,
|
||||
[&](int64_t val) { ss << "static_cast<int64_t>(" << val << ")"; },
|
||||
", ");
|
||||
attrDefs.push_back(llvm::formatv(defFmt, llvm::formatv(typeFmt, size),
|
||||
|
@ -1088,11 +1088,11 @@ if ({1}Iter != attrs.end()) {{
|
|||
if (expression.scalarFn->attrName) {
|
||||
if (llvm::none_of(args, [&](LinalgOperandDef &arg) {
|
||||
return isFunctionAttribute(arg.kind) &&
|
||||
arg.name == expression.scalarFn->attrName.getValue();
|
||||
arg.name == expression.scalarFn->attrName.value();
|
||||
})) {
|
||||
emitError(genContext.getLoc())
|
||||
<< "missing function attribute "
|
||||
<< expression.scalarFn->attrName.getValue();
|
||||
<< expression.scalarFn->attrName.value();
|
||||
}
|
||||
funcType = llvm::formatv("{0}Val", *expression.scalarFn->attrName);
|
||||
}
|
||||
|
@ -1103,15 +1103,15 @@ if ({1}Iter != attrs.end()) {{
|
|||
if (expression.scalarFn->kind == ScalarFnKind::Type) {
|
||||
assert(expression.scalarFn->typeVar.has_value());
|
||||
Optional<std::string> typeCppValue =
|
||||
findTypeValue(expression.scalarFn->typeVar.getValue(), args);
|
||||
findTypeValue(expression.scalarFn->typeVar.value(), args);
|
||||
if (!typeCppValue) {
|
||||
emitError(genContext.getLoc())
|
||||
<< "type variable " << expression.scalarFn->typeVar.getValue()
|
||||
<< "type variable " << expression.scalarFn->typeVar.value()
|
||||
<< ", used in a type conversion, must map to a predefined or "
|
||||
<< "an argument type but it does not";
|
||||
return None;
|
||||
}
|
||||
operandCppValues.push_back(typeCppValue.getValue());
|
||||
operandCppValues.push_back(typeCppValue.value());
|
||||
}
|
||||
|
||||
// Collect the scalar operands.
|
||||
|
|
Loading…
Reference in New Issue