diff --git a/flang/lib/Optimizer/Transforms/AffineDemotion.cpp b/flang/lib/Optimizer/Transforms/AffineDemotion.cpp index 15d8a4ece69f8..96e0853b24b5e 100644 --- a/flang/lib/Optimizer/Transforms/AffineDemotion.cpp +++ b/flang/lib/Optimizer/Transforms/AffineDemotion.cpp @@ -50,7 +50,7 @@ class AffineLoadConversion : public OpConversionPattern { LogicalResult matchAndRewrite(mlir::AffineLoadOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - SmallVector indices(adaptor.indices()); + SmallVector indices(adaptor.getIndices()); auto maybeExpandedMap = expandAffineMap(rewriter, op.getLoc(), op.getAffineMap(), indices); if (!maybeExpandedMap) @@ -58,7 +58,7 @@ class AffineLoadConversion : public OpConversionPattern { auto coorOp = rewriter.create( op.getLoc(), fir::ReferenceType::get(op.getResult().getType()), - adaptor.memref(), *maybeExpandedMap); + adaptor.getMemref(), *maybeExpandedMap); rewriter.replaceOpWithNewOp(op, coorOp.getResult()); return success(); @@ -72,7 +72,7 @@ class AffineStoreConversion : public OpConversionPattern { LogicalResult matchAndRewrite(mlir::AffineStoreOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - SmallVector indices(op.indices()); + SmallVector indices(op.getIndices()); auto maybeExpandedMap = expandAffineMap(rewriter, op.getLoc(), op.getAffineMap(), indices); if (!maybeExpandedMap) @@ -80,8 +80,8 @@ class AffineStoreConversion : public OpConversionPattern { auto coorOp = rewriter.create( op.getLoc(), fir::ReferenceType::get(op.getValueToStore().getType()), - adaptor.memref(), *maybeExpandedMap); - rewriter.replaceOpWithNewOp(op, adaptor.value(), + adaptor.getMemref(), *maybeExpandedMap); + rewriter.replaceOpWithNewOp(op, adaptor.getValue(), coorOp.getResult()); return success(); } diff --git a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp index b6f5c901adef7..b50b766d98b4b 100644 --- a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp +++ b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp @@ -34,8 +34,8 @@ struct RawBufferOpLowering : public ConvertOpToLLVMPattern { matchAndRewrite(GpuOp gpuOp, typename GpuOp::Adaptor adaptor, ConversionPatternRewriter &rewriter) const override { Location loc = gpuOp.getLoc(); - Value memref = adaptor.memref(); - Value unconvertedMemref = gpuOp.memref(); + Value memref = adaptor.getMemref(); + Value unconvertedMemref = gpuOp.getMemref(); MemRefType memrefType = unconvertedMemref.getType().cast(); Value storeData = adaptor.getODSOperands(0)[0]; @@ -163,9 +163,9 @@ struct RawBufferOpLowering : public ConvertOpToLLVMPattern { // swizzles) RDNA only // bits 30-31: Type (must be 0) uint32_t word3 = (7 << 12) | (4 << 15); - if (adaptor.targetIsRDNA()) { + if (adaptor.getTargetIsRDNA()) { word3 |= (1 << 24); - uint32_t oob = adaptor.boundsCheck() ? 1 : 2; + uint32_t oob = adaptor.getBoundsCheck() ? 1 : 2; word3 |= (oob << 28); } Value word3Const = createI32Constant(rewriter, loc, word3); @@ -176,7 +176,7 @@ struct RawBufferOpLowering : public ConvertOpToLLVMPattern { // Indexing (voffset) Value voffset; - for (auto &pair : llvm::enumerate(adaptor.indices())) { + for (auto &pair : llvm::enumerate(adaptor.getIndices())) { size_t i = pair.index(); Value index = pair.value(); Value strideOp; @@ -191,8 +191,8 @@ struct RawBufferOpLowering : public ConvertOpToLLVMPattern { voffset = voffset ? rewriter.create(loc, voffset, index) : index; } - if (adaptor.indexOffset().hasValue()) { - int32_t indexOffset = *gpuOp.indexOffset() * elementByteWidth; + if (adaptor.getIndexOffset().hasValue()) { + int32_t indexOffset = *gpuOp.getIndexOffset() * elementByteWidth; Value extraOffsetConst = createI32Constant(rewriter, loc, indexOffset); voffset = voffset ? rewriter.create(loc, voffset, extraOffsetConst) @@ -200,7 +200,7 @@ struct RawBufferOpLowering : public ConvertOpToLLVMPattern { } args.push_back(voffset); - Value sgprOffset = adaptor.sgprOffset(); + Value sgprOffset = adaptor.getSgprOffset(); if (!sgprOffset) sgprOffset = createI32Constant(rewriter, loc, 0); if (ShapedType::isDynamicStrideOrOffset(offset)) diff --git a/mlir/lib/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.cpp b/mlir/lib/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.cpp index e8c74c98319ff..a2f5641a0f5e2 100644 --- a/mlir/lib/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.cpp +++ b/mlir/lib/Conversion/ArmNeon2dToIntr/ArmNeon2dToIntr.cpp @@ -28,19 +28,19 @@ class Sdot2dLoweringPattern : public OpRewritePattern { /// arm.neon.intr.sdot LogicalResult matchAndRewrite(Sdot2dOp op, PatternRewriter &rewriter) const override { - Type elemType = op.b().getType().cast().getElementType(); - int length = op.b().getType().cast().getShape()[0] * + Type elemType = op.getB().getType().cast().getElementType(); + int length = op.getB().getType().cast().getShape()[0] * Sdot2dOp::kReductionSize; VectorType flattenedVectorType = VectorType::get({length}, elemType); - Value b2d = op.b(); - Value c2d = op.c(); + Value b2d = op.getB(); + Value c2d = op.getC(); Location loc = op.getLoc(); Value b1d = rewriter.create(loc, flattenedVectorType, b2d); Value c1d = rewriter.create(loc, flattenedVectorType, c2d); - Value newOp = - rewriter.create(loc, op.res().getType(), op.a(), b1d, c1d); + Value newOp = rewriter.create(loc, op.getRes().getType(), op.getA(), + b1d, c1d); rewriter.replaceOp(op, {newOp}); return success(); } diff --git a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp index 41388a60eee2e..f980e52c3965b 100644 --- a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp +++ b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp @@ -221,14 +221,15 @@ struct MmaLdMatrixOpToNVVM : public ConvertOpToLLVMPattern { ldMatrixResultType = rewriter.getI32Type(); } - auto srcMemrefType = op.srcMemref().getType().cast(); - Value srcPtr = getStridedElementPtr(loc, srcMemrefType, adaptor.srcMemref(), - adaptor.indices(), rewriter); + auto srcMemrefType = op.getSrcMemref().getType().cast(); + Value srcPtr = + getStridedElementPtr(loc, srcMemrefType, adaptor.getSrcMemref(), + adaptor.getIndices(), rewriter); Value ldMatrixResult = rewriter.create( loc, ldMatrixResultType, srcPtr, - /*num=*/op.numTiles(), - /*layout=*/op.transpose() ? NVVM::MMALayout::col - : NVVM::MMALayout::row); + /*num=*/op.getNumTiles(), + /*layout=*/op.getTranspose() ? NVVM::MMALayout::col + : NVVM::MMALayout::row); // The ldmatrix operation returns either a single i32 value or a struct of // i32 values. Here we unpack those values and cast them back to their @@ -262,12 +263,12 @@ struct MmaSyncOptoNVVM : public ConvertOpToLLVMPattern { Location loc = op->getLoc(); // Get the shapes of the MMAMatrix type being used. The shapes will // choose which intrinsic this op will be lowered to. - auto aType = op.matrixA().getType().cast(); - auto cType = op.matrixC().getType().cast(); + auto aType = op.getMatrixA().getType().cast(); + auto cType = op.getMatrixC().getType().cast(); - int64_t m = op.mmaShape()[0].cast().getInt(); - int64_t n = op.mmaShape()[1].cast().getInt(); - int64_t k = op.mmaShape()[2].cast().getInt(); + int64_t m = op.getMmaShape()[0].cast().getInt(); + int64_t n = op.getMmaShape()[1].cast().getInt(); + int64_t k = op.getMmaShape()[2].cast().getInt(); std::array gemmShape{m, n, k}; NVVM::MMATypes ptxTypeA; @@ -302,11 +303,11 @@ struct MmaSyncOptoNVVM : public ConvertOpToLLVMPattern { } SmallVector matA = - unpackOperandVector(rewriter, loc, adaptor.matrixA(), ptxTypeA); + unpackOperandVector(rewriter, loc, adaptor.getMatrixA(), ptxTypeA); SmallVector matB = - unpackOperandVector(rewriter, loc, adaptor.matrixB(), ptxTypeB); + unpackOperandVector(rewriter, loc, adaptor.getMatrixB(), ptxTypeB); SmallVector matC = - unpackOperandVector(rewriter, loc, adaptor.matrixC(), *ptxTypeC); + unpackOperandVector(rewriter, loc, adaptor.getMatrixC(), *ptxTypeC); Type desiredRetTy = typeConverter->convertType(op->getResultTypes()[0]); Type intrinsicResTy = inferIntrinsicResultType( @@ -359,18 +360,18 @@ struct NVGPUAsyncCopyLowering matchAndRewrite(nvgpu::DeviceAsyncCopyOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { Location loc = op->getLoc(); - auto dstMemrefType = op.dst().getType().cast(); - Value dstPtr = getStridedElementPtr(loc, dstMemrefType, adaptor.dst(), - adaptor.dstIndices(), rewriter); + auto dstMemrefType = op.getDst().getType().cast(); + Value dstPtr = getStridedElementPtr(loc, dstMemrefType, adaptor.getDst(), + adaptor.getDstIndices(), rewriter); auto i8Ty = IntegerType::get(op.getContext(), 8); auto dstPointerType = LLVM::LLVMPointerType::get(i8Ty, dstMemrefType.getMemorySpaceAsInt()); dstPtr = rewriter.create(loc, dstPointerType, dstPtr); - auto srcMemrefType = op.src().getType().cast(); + auto srcMemrefType = op.getSrc().getType().cast(); - Value scrPtr = getStridedElementPtr(loc, srcMemrefType, adaptor.src(), - adaptor.srcIndices(), rewriter); + Value scrPtr = getStridedElementPtr(loc, srcMemrefType, adaptor.getSrc(), + adaptor.getSrcIndices(), rewriter); auto srcPointerType = LLVM::LLVMPointerType::get(i8Ty, srcMemrefType.getMemorySpaceAsInt()); scrPtr = rewriter.create(loc, srcPointerType, scrPtr); @@ -379,12 +380,13 @@ struct NVGPUAsyncCopyLowering i8Ty, NVVM::NVVMMemorySpace::kGlobalMemorySpace); scrPtr = rewriter.create(loc, srcPointerGlobalType, scrPtr); - int64_t numElements = adaptor.numElements().getZExtValue(); + int64_t numElements = adaptor.getNumElements().getZExtValue(); int64_t sizeInBytes = (dstMemrefType.getElementTypeBitWidth() * numElements) / 8; // bypass L1 is only supported for byte sizes of 16, we drop the hint // otherwise. - UnitAttr bypassL1 = sizeInBytes == 16 ? adaptor.bypassL1Attr() : UnitAttr(); + UnitAttr bypassL1 = + sizeInBytes == 16 ? adaptor.getBypassL1Attr() : UnitAttr(); rewriter.create( loc, dstPtr, scrPtr, rewriter.getI32IntegerAttr(sizeInBytes), bypassL1); @@ -424,7 +426,7 @@ struct NVGPUAsyncWaitLowering matchAndRewrite(nvgpu::DeviceAsyncWaitOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { // If numGroup is not present pick 0 as a conservative correct value. - int32_t numGroups = adaptor.numGroups() ? *adaptor.numGroups() : 0; + int32_t numGroups = adaptor.getNumGroups() ? *adaptor.getNumGroups() : 0; rewriter.create(op.getLoc(), numGroups); rewriter.eraseOp(op); return success(); diff --git a/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp b/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp index 6e74a90bc03d9..fc5bf877c9a91 100644 --- a/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp +++ b/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp @@ -106,7 +106,7 @@ static Value getOrEmitUpperBound(AffineForOp forOp, OpBuilder &builder) { // rewriting infrastructure. static LogicalResult checkAffineLoopNestMappableImpl(AffineForOp forOp, unsigned numDims) { - Region &limit = forOp.region(); + Region &limit = forOp.getRegion(); for (unsigned i = 0, e = numDims; i < e; ++i) { Operation *nested = &forOp.getBody()->front(); if (!areValuesDefinedAbove(getLowerBoundOperands(forOp), limit) || @@ -320,7 +320,7 @@ static Value deriveStaticUpperBound(Value upperBound, } if (auto minOp = upperBound.getDefiningOp()) { - for (const AffineExpr &result : minOp.map().getResults()) { + for (const AffineExpr &result : minOp.getMap().getResults()) { if (auto constExpr = result.dyn_cast()) { return rewriter.create(minOp.getLoc(), constExpr.getValue()); diff --git a/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRV.cpp b/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRV.cpp index 04795049c1d68..44ba0d0adaab4 100644 --- a/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRV.cpp +++ b/mlir/lib/Conversion/TensorToSPIRV/TensorToSPIRV.cpp @@ -45,7 +45,7 @@ class TensorExtractPattern final LogicalResult matchAndRewrite(tensor::ExtractOp extractOp, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - TensorType tensorType = extractOp.tensor().getType().cast(); + TensorType tensorType = extractOp.getTensor().getType().cast(); if (!tensorType.hasStaticShape()) return rewriter.notifyMatchFailure(extractOp, "non-static tensor"); @@ -63,14 +63,14 @@ class TensorExtractPattern final strides[i] = strides[i + 1] * tensorType.getDimSize(i + 1); } - Type varType = spirv::PointerType::get(adaptor.tensor().getType(), + Type varType = spirv::PointerType::get(adaptor.getTensor().getType(), spirv::StorageClass::Function); spirv::VariableOp varOp; - if (adaptor.tensor().getDefiningOp()) { + if (adaptor.getTensor().getDefiningOp()) { varOp = rewriter.create( loc, varType, spirv::StorageClass::Function, - /*initializer=*/adaptor.tensor()); + /*initializer=*/adaptor.getTensor()); } else { // Need to store the value to the local variable. It's questionable // whether we want to support such case though. @@ -80,7 +80,7 @@ class TensorExtractPattern final auto &typeConverter = *getTypeConverter(); auto indexType = typeConverter.getIndexType(); - Value index = spirv::linearizeIndex(adaptor.indices(), strides, + Value index = spirv::linearizeIndex(adaptor.getIndices(), strides, /*offset=*/0, indexType, loc, rewriter); auto acOp = rewriter.create(loc, varOp, index); diff --git a/mlir/lib/Dialect/AMDGPU/IR/AMDGPUDialect.cpp b/mlir/lib/Dialect/AMDGPU/IR/AMDGPUDialect.cpp index 08dc1351c4d23..68639da21428b 100644 --- a/mlir/lib/Dialect/AMDGPU/IR/AMDGPUDialect.cpp +++ b/mlir/lib/Dialect/AMDGPU/IR/AMDGPUDialect.cpp @@ -31,14 +31,14 @@ void amdgpu::AMDGPUDialect::initialize() { //===----------------------------------------------------------------------===// template static LogicalResult verifyRawBufferOp(T &op) { - MemRefType bufferType = op.memref().getType().template cast(); + MemRefType bufferType = op.getMemref().getType().template cast(); if (bufferType.getMemorySpaceAsInt() != 0) return op.emitOpError( "Buffer ops must operate on a memref in global memory"); if (!bufferType.hasRank()) return op.emitOpError( "Cannot meaningfully buffer_store to an unranked memref"); - if (static_cast(op.indices().size()) != bufferType.getRank()) + if (static_cast(op.getIndices().size()) != bufferType.getRank()) return op.emitOpError("Expected " + Twine(bufferType.getRank()) + " indices to memref"); return success(); diff --git a/mlir/lib/Dialect/AMX/IR/AMXDialect.cpp b/mlir/lib/Dialect/AMX/IR/AMXDialect.cpp index 9ea96791cef4b..f0e434407c8a2 100644 --- a/mlir/lib/Dialect/AMX/IR/AMXDialect.cpp +++ b/mlir/lib/Dialect/AMX/IR/AMXDialect.cpp @@ -58,14 +58,14 @@ LogicalResult amx::TileZeroOp::verify() { LogicalResult amx::TileLoadOp::verify() { unsigned rank = getMemRefType().getRank(); - if (indices().size() != rank) + if (getIndices().size() != rank) return emitOpError("requires ") << rank << " indices"; return verifyTileSize(*this, getVectorType()); } LogicalResult amx::TileStoreOp::verify() { unsigned rank = getMemRefType().getRank(); - if (indices().size() != rank) + if (getIndices().size() != rank) return emitOpError("requires ") << rank << " indices"; return verifyTileSize(*this, getVectorType()); } diff --git a/mlir/lib/Dialect/AMX/Transforms/LegalizeForLLVMExport.cpp b/mlir/lib/Dialect/AMX/Transforms/LegalizeForLLVMExport.cpp index e6949fa862e46..c19f8f182a923 100644 --- a/mlir/lib/Dialect/AMX/Transforms/LegalizeForLLVMExport.cpp +++ b/mlir/lib/Dialect/AMX/Transforms/LegalizeForLLVMExport.cpp @@ -112,10 +112,10 @@ struct TileLoadConversion : public ConvertOpToLLVMPattern { if (failed(verifyStride(mType))) return failure(); Value stride = getStride(rewriter, *getTypeConverter(), mType, - adaptor.base(), op.getLoc()); + adaptor.getBase(), op.getLoc()); // Replace operation with intrinsic. - Value ptr = getStridedElementPtr(op.getLoc(), mType, adaptor.base(), - adaptor.indices(), rewriter); + Value ptr = getStridedElementPtr(op.getLoc(), mType, adaptor.getBase(), + adaptor.getIndices(), rewriter); ptr = castPtr(rewriter, op.getLoc(), ptr); Type resType = typeConverter->convertType(vType); rewriter.replaceOpWithNewOp( @@ -139,13 +139,13 @@ struct TileStoreConversion : public ConvertOpToLLVMPattern { if (failed(verifyStride(mType))) return failure(); Value stride = getStride(rewriter, *getTypeConverter(), mType, - adaptor.base(), op.getLoc()); + adaptor.getBase(), op.getLoc()); // Replace operation with intrinsic. - Value ptr = getStridedElementPtr(op.getLoc(), mType, adaptor.base(), - adaptor.indices(), rewriter); + Value ptr = getStridedElementPtr(op.getLoc(), mType, adaptor.getBase(), + adaptor.getIndices(), rewriter); ptr = castPtr(rewriter, op.getLoc(), ptr); rewriter.replaceOpWithNewOp( - op, tsz.first, tsz.second, ptr, stride, adaptor.val()); + op, tsz.first, tsz.second, ptr, stride, adaptor.getVal()); return success(); } }; @@ -166,8 +166,8 @@ struct TileMulFConversion : public ConvertOpToLLVMPattern { // Replace operation with intrinsic. Type resType = typeConverter->convertType(cType); rewriter.replaceOpWithNewOp( - op, resType, tsza.first, tszb.second, tsza.second, adaptor.acc(), - adaptor.lhs(), adaptor.rhs()); + op, resType, tsza.first, tszb.second, tsza.second, adaptor.getAcc(), + adaptor.getLhs(), adaptor.getRhs()); return success(); } }; @@ -187,24 +187,24 @@ struct TileMulIConversion : public ConvertOpToLLVMPattern { getTileSizes(rewriter, *getTypeConverter(), bType, op.getLoc()); // Replace operation with intrinsic. Type resType = typeConverter->convertType(cType); - bool zexta = op.isZextLhs(); - bool zextb = op.isZextRhs(); + bool zexta = op.getIsZextLhs(); + bool zextb = op.getIsZextRhs(); if (zexta && zextb) rewriter.replaceOpWithNewOp( - op, resType, tsza.first, tszb.second, tsza.second, adaptor.acc(), - adaptor.lhs(), adaptor.rhs()); + op, resType, tsza.first, tszb.second, tsza.second, adaptor.getAcc(), + adaptor.getLhs(), adaptor.getRhs()); else if (zexta && !zextb) rewriter.replaceOpWithNewOp( - op, resType, tsza.first, tszb.second, tsza.second, adaptor.acc(), - adaptor.lhs(), adaptor.rhs()); + op, resType, tsza.first, tszb.second, tsza.second, adaptor.getAcc(), + adaptor.getLhs(), adaptor.getRhs()); else if (!zexta && zextb) rewriter.replaceOpWithNewOp( - op, resType, tsza.first, tszb.second, tsza.second, adaptor.acc(), - adaptor.lhs(), adaptor.rhs()); + op, resType, tsza.first, tszb.second, tsza.second, adaptor.getAcc(), + adaptor.getLhs(), adaptor.getRhs()); else rewriter.replaceOpWithNewOp( - op, resType, tsza.first, tszb.second, tsza.second, adaptor.acc(), - adaptor.lhs(), adaptor.rhs()); + op, resType, tsza.first, tszb.second, tsza.second, adaptor.getAcc(), + adaptor.getLhs(), adaptor.getRhs()); return success(); } }; diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp index d217f6bfe02c5..8883c4aaf6f93 100644 --- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp +++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp @@ -65,9 +65,9 @@ ParseResult VoteBallotOp::parse(OpAsmParser &parser, OperationState &result) { void VoteBallotOp::print(OpAsmPrinter &p) { printNVVMIntrinsicOp(p, *this); } LogicalResult CpAsyncOp::verify() { - if (size() != 4 && size() != 8 && size() != 16) + if (getSize() != 4 && getSize() != 8 && getSize() != 16) return emitError("expected byte size to be either 4, 8 or 16."); - if (bypass_l1() && size() != 16) + if (getBypassL1() && getSize() != 16) return emitError("bypass l1 is only support for 16 bytes copy."); return success(); } @@ -140,8 +140,8 @@ void MmaOp::print(OpAsmPrinter &p) { }; std::array frags{ - OperandFragment("A", multiplicandAPtxTypeAttrName()), - OperandFragment("B", multiplicandBPtxTypeAttrName()), + OperandFragment("A", getMultiplicandAPtxTypeAttrName()), + OperandFragment("B", getMultiplicandBPtxTypeAttrName()), OperandFragment("C", "")}; SmallVector ignoreAttrNames{ mlir::NVVM::MmaOp::getOperandSegmentSizeAttr()}; @@ -184,7 +184,7 @@ void MmaOp::print(OpAsmPrinter &p) { frags[2].regs[0].getType()}, p); p << ")"; - p.printArrowTypeList(TypeRange{this->res().getType()}); + p.printArrowTypeList(TypeRange{this->getRes().getType()}); } void MmaOp::build(OpBuilder &builder, OperationState &result, Type resultType, @@ -355,8 +355,8 @@ LogicalResult MmaOp::verify() { auto s32x2StructTy = LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty}); - std::array mmaShape{shapeAttr().getM(), shapeAttr().getN(), - shapeAttr().getK()}; + std::array mmaShape{getShapeAttr().getM(), getShapeAttr().getN(), + getShapeAttr().getK()}; // These variables define the set of allowed data types for matrices A, B, C, // and result. @@ -373,7 +373,7 @@ LogicalResult MmaOp::verify() { if (mmaShape[0] == 16) { int64_t kFactor; Type multiplicandFragType; - switch (multiplicandAPtxType().getValue()) { + switch (getMultiplicandAPtxType().getValue()) { case MMATypes::tf32: kFactor = 4; multiplicandFragType = i32Ty; @@ -400,10 +400,10 @@ LogicalResult MmaOp::verify() { break; default: return emitError("invalid shape or multiplicand type: " + - stringifyEnum(multiplicandAPtxType().getValue())); + stringifyEnum(getMultiplicandAPtxType().getValue())); } - if (isIntegerPtxType(multiplicandAPtxType().getValue())) { + if (isIntegerPtxType(getMultiplicandAPtxType().getValue())) { expectedResult.push_back(s32x4StructTy); expectedC.emplace_back(4, i32Ty); multiplicandFragType = i32Ty; @@ -422,7 +422,7 @@ LogicalResult MmaOp::verify() { // In the M=8 case, there is only 1 possible case per data type. if (mmaShape[0] == 8) { - if (multiplicandAPtxType().getValue() == MMATypes::f16) { + if (getMultiplicandAPtxType().getValue() == MMATypes::f16) { expectedA.emplace_back(2, f16x2Ty); expectedB.emplace_back(2, f16x2Ty); expectedResult.push_back(f16x2x4StructTy); @@ -431,7 +431,7 @@ LogicalResult MmaOp::verify() { expectedC.emplace_back(8, f32Ty); allowedShapes.push_back({8, 8, 4}); } - if (multiplicandAPtxType().getValue() == MMATypes::f64) { + if (getMultiplicandAPtxType().getValue() == MMATypes::f64) { Type f64Ty = Float64Type::get(context); expectedA.emplace_back(1, f64Ty); expectedB.emplace_back(1, f64Ty); @@ -441,16 +441,16 @@ LogicalResult MmaOp::verify() { context, SmallVector(2, f64Ty))); allowedShapes.push_back({8, 8, 4}); } - if (isIntegerPtxType(multiplicandAPtxType().getValue())) { + if (isIntegerPtxType(getMultiplicandAPtxType().getValue())) { expectedA.push_back({i32Ty}); expectedB.push_back({i32Ty}); expectedC.push_back({i32Ty, i32Ty}); expectedResult.push_back(s32x2StructTy); - if (isInt4PtxType(multiplicandAPtxType().getValue())) + if (isInt4PtxType(getMultiplicandAPtxType().getValue())) allowedShapes.push_back({8, 8, 32}); - if (isInt8PtxType(multiplicandAPtxType().getValue())) + if (isInt8PtxType(getMultiplicandAPtxType().getValue())) allowedShapes.push_back({8, 8, 16}); - if (multiplicandAPtxType().getValue() == MMATypes::b1) + if (getMultiplicandAPtxType().getValue() == MMATypes::b1) allowedShapes.push_back({8, 8, 128}); } } @@ -506,17 +506,19 @@ LogicalResult MmaOp::verify() { } // Ensure that binary MMA variants have a b1 MMA operation defined. - if (multiplicandAPtxType() == MMATypes::b1 && !b1Op().hasValue()) { - return emitOpError("op requires " + b1OpAttrName().strref() + " attribute"); + if (getMultiplicandAPtxType() == MMATypes::b1 && !getB1Op().hasValue()) { + return emitOpError("op requires " + getB1OpAttrName().strref() + + " attribute"); } // Ensure int4/int8 MMA variants specify the accum overflow behavior // attribute. - if (isInt4PtxType(*multiplicandAPtxType()) || - isInt8PtxType(*multiplicandAPtxType())) { - if (!intOverflowBehavior().hasValue()) + if (isInt4PtxType(*getMultiplicandAPtxType()) || + isInt8PtxType(*getMultiplicandAPtxType())) { + if (!getIntOverflowBehavior().hasValue()) return emitOpError("op requires " + - intOverflowBehaviorAttrName().strref() + " attribute"); + getIntOverflowBehaviorAttrName().strref() + + " attribute"); } return success(); @@ -561,16 +563,16 @@ std::pair NVVM::inferMMAType(NVVM::MMATypes type, LogicalResult NVVM::WMMALoadOp::verify() { unsigned addressSpace = - ptr().getType().cast().getAddressSpace(); + getPtr().getType().cast().getAddressSpace(); if (addressSpace != 0 && addressSpace != 1 && addressSpace != 3) return emitOpError("expected source pointer in memory " "space 0, 1, 3"); - if (NVVM::WMMALoadOp::getIntrinsicID(m(), n(), k(), layout(), eltype(), - frag()) == 0) + if (NVVM::WMMALoadOp::getIntrinsicID(getM(), getN(), getK(), getLayout(), + getEltype(), getFrag()) == 0) return emitOpError() << "invalid attribute combination"; std::pair typeInfo = - inferMMAType(eltype(), frag(), getContext()); + inferMMAType(getEltype(), getFrag(), getContext()); Type dstType = LLVM::LLVMStructType::getLiteral( getContext(), SmallVector(typeInfo.second, typeInfo.first)); if (getType() != dstType) @@ -581,18 +583,19 @@ LogicalResult NVVM::WMMALoadOp::verify() { LogicalResult NVVM::WMMAStoreOp::verify() { unsigned addressSpace = - ptr().getType().cast().getAddressSpace(); + getPtr().getType().cast().getAddressSpace(); if (addressSpace != 0 && addressSpace != 1 && addressSpace != 3) return emitOpError("expected operands to be a source pointer in memory " "space 0, 1, 3"); - if (NVVM::WMMAStoreOp::getIntrinsicID(m(), n(), k(), layout(), eltype()) == 0) + if (NVVM::WMMAStoreOp::getIntrinsicID(getM(), getN(), getK(), getLayout(), + getEltype()) == 0) return emitOpError() << "invalid attribute combination"; std::pair typeInfo = - inferMMAType(eltype(), NVVM::MMAFrag::c, getContext()); - if (args().size() != typeInfo.second) + inferMMAType(getEltype(), NVVM::MMAFrag::c, getContext()); + if (getArgs().size() != typeInfo.second) return emitOpError() << "expected " << typeInfo.second << " data operands"; - if (llvm::any_of(args(), [&typeInfo](Value operands) { + if (llvm::any_of(getArgs(), [&typeInfo](Value operands) { return operands.getType() != typeInfo.first; })) return emitOpError() << "expected data operands of type " << typeInfo.first; @@ -600,24 +603,25 @@ LogicalResult NVVM::WMMAStoreOp::verify() { } LogicalResult NVVM::WMMAMmaOp::verify() { - if (NVVM::WMMAMmaOp::getIntrinsicID(m(), n(), k(), layoutA(), layoutB(), - eltypeA(), eltypeB()) == 0) + if (NVVM::WMMAMmaOp::getIntrinsicID(getM(), getN(), getK(), getLayoutA(), + getLayoutB(), getEltypeA(), + getEltypeB()) == 0) return emitOpError() << "invalid attribute combination"; std::pair typeInfoA = - inferMMAType(eltypeA(), NVVM::MMAFrag::a, getContext()); + inferMMAType(getEltypeA(), NVVM::MMAFrag::a, getContext()); std::pair typeInfoB = - inferMMAType(eltypeA(), NVVM::MMAFrag::b, getContext()); + inferMMAType(getEltypeA(), NVVM::MMAFrag::b, getContext()); std::pair typeInfoC = - inferMMAType(eltypeB(), NVVM::MMAFrag::c, getContext()); + inferMMAType(getEltypeB(), NVVM::MMAFrag::c, getContext()); SmallVector arguments; arguments.append(typeInfoA.second, typeInfoA.first); arguments.append(typeInfoB.second, typeInfoB.first); arguments.append(typeInfoC.second, typeInfoC.first); unsigned numArgs = arguments.size(); - if (args().size() != numArgs) + if (getArgs().size() != numArgs) return emitOpError() << "expected " << numArgs << " arguments"; for (unsigned i = 0; i < numArgs; i++) { - if (args()[i].getType() != arguments[i]) + if (getArgs()[i].getType() != arguments[i]) return emitOpError() << "expected argument " << i << " to be of type " << arguments[i]; } @@ -631,22 +635,22 @@ LogicalResult NVVM::WMMAMmaOp::verify() { LogicalResult NVVM::LdMatrixOp::verify() { unsigned addressSpace = - ptr().getType().cast().getAddressSpace(); + getPtr().getType().cast().getAddressSpace(); if (addressSpace != 3) return emitOpError("expected source pointer in memory space 3"); - if (num() != 1 && num() != 2 && num() != 4) + if (getNum() != 1 && getNum() != 2 && getNum() != 4) return emitOpError("expected num attribute to be 1, 2 or 4"); Type i32 = IntegerType::get(getContext(), 32); - if (num() == 1 && getType() != i32) + if (getNum() == 1 && getType() != i32) return emitOpError("expected destination type is i32"); - if (num() == 2 || num() == 4) { + if (getNum() == 2 || getNum() == 4) { Type dstType = LLVM::LLVMStructType::getLiteral( - getContext(), SmallVector(num(), i32)); + getContext(), SmallVector(getNum(), i32)); if (getType() != dstType) return emitOpError("expected destination type is a structure of ") - << num() << " elements of type i32"; + << getNum() << " elements of type i32"; } return success(); } diff --git a/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp index 87c574eb8f672..338e71517f4cc 100644 --- a/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp +++ b/mlir/lib/Dialect/LLVMIR/IR/ROCDLDialect.cpp @@ -81,7 +81,7 @@ ParseResult MubufStoreOp::parse(OpAsmParser &parser, OperationState &result) { } void MubufStoreOp::print(OpAsmPrinter &p) { - p << " " << getOperands() << " : " << vdata().getType(); + p << " " << getOperands() << " : " << getVdata().getType(); } // ::= @@ -103,7 +103,7 @@ ParseResult RawBufferLoadOp::parse(OpAsmParser &parser, } void RawBufferLoadOp::print(OpAsmPrinter &p) { - p << " " << getOperands() << " : " << res().getType(); + p << " " << getOperands() << " : " << getRes().getType(); } // ::= @@ -127,7 +127,7 @@ ParseResult RawBufferStoreOp::parse(OpAsmParser &parser, } void RawBufferStoreOp::print(OpAsmPrinter &p) { - p << " " << getOperands() << " : " << vdata().getType(); + p << " " << getOperands() << " : " << getVdata().getType(); } // ::= @@ -151,7 +151,7 @@ ParseResult RawBufferAtomicFAddOp::parse(OpAsmParser &parser, } void RawBufferAtomicFAddOp::print(mlir::OpAsmPrinter &p) { - p << " " << getOperands() << " : " << vdata().getType(); + p << " " << getOperands() << " : " << getVdata().getType(); } //===----------------------------------------------------------------------===// diff --git a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp index 1a781ef364ffe..c31a168cd2103 100644 --- a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp +++ b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp @@ -66,8 +66,8 @@ static bool isLastMemrefDimUnitStride(MemRefType type) { } LogicalResult DeviceAsyncCopyOp::verify() { - auto srcMemref = src().getType().cast(); - auto dstMemref = dst().getType().cast(); + auto srcMemref = getSrc().getType().cast(); + auto dstMemref = getDst().getType().cast(); unsigned workgroupAddressSpace = gpu::GPUDialect::getWorkgroupAddressSpace(); if (!isLastMemrefDimUnitStride(srcMemref)) return emitError("source memref most minor dim must have unit stride"); @@ -78,12 +78,13 @@ LogicalResult DeviceAsyncCopyOp::verify() { << workgroupAddressSpace; if (dstMemref.getElementType() != srcMemref.getElementType()) return emitError("source and destination must have the same element type"); - if (size_t(srcMemref.getRank()) != srcIndices().size()) + if (size_t(srcMemref.getRank()) != getSrcIndices().size()) return emitOpError() << "expected " << srcMemref.getRank() - << " source indices, got " << srcIndices().size(); - if (size_t(dstMemref.getRank()) != dstIndices().size()) + << " source indices, got " << getSrcIndices().size(); + if (size_t(dstMemref.getRank()) != getDstIndices().size()) return emitOpError() << "expected " << dstMemref.getRank() - << " destination indices, got " << dstIndices().size(); + << " destination indices, got " + << getDstIndices().size(); return success(); } diff --git a/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp b/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp index 3d01e2ee0998e..1760bde459488 100644 --- a/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp +++ b/mlir/lib/Dialect/NVGPU/Transforms/OptimizeSharedMemory.cpp @@ -101,9 +101,9 @@ static void transformIndices(OpBuilder &builder, Location loc, Operation::operand_range getIndices(Operation *op) { if (auto ldmatrixOp = dyn_cast(op)) - return ldmatrixOp.indices(); + return ldmatrixOp.getIndices(); if (auto copyOp = dyn_cast(op)) - return copyOp.dstIndices(); + return copyOp.getDstIndices(); if (auto loadOp = dyn_cast(op)) return loadOp.indices(); if (auto storeOp = dyn_cast(op)) @@ -117,9 +117,9 @@ Operation::operand_range getIndices(Operation *op) { void setIndices(Operation *op, ArrayRef indices) { if (auto ldmatrixOp = dyn_cast(op)) - return ldmatrixOp.indicesMutable().assign(indices); + return ldmatrixOp.getIndicesMutable().assign(indices); if (auto copyOp = dyn_cast(op)) - return copyOp.dstIndicesMutable().assign(indices); + return copyOp.getDstIndicesMutable().assign(indices); if (auto loadOp = dyn_cast(op)) return loadOp.indicesMutable().assign(indices); if (auto storeOp = dyn_cast(op)) diff --git a/mlir/lib/Dialect/SCF/Transforms/BufferizableOpInterfaceImpl.cpp b/mlir/lib/Dialect/SCF/Transforms/BufferizableOpInterfaceImpl.cpp index 4bff6b56e240c..71ca88dffdd9b 100644 --- a/mlir/lib/Dialect/SCF/Transforms/BufferizableOpInterfaceImpl.cpp +++ b/mlir/lib/Dialect/SCF/Transforms/BufferizableOpInterfaceImpl.cpp @@ -1036,7 +1036,7 @@ static bool areEquivalentExtractSliceOps(const AnalysisState &state, if (!st || !sti) return false; if (st != sti && - !state.areEquivalentBufferizedValues(st.source(), sti.getDest())) + !state.areEquivalentBufferizedValues(st.getSource(), sti.getDest())) return false; if (!sameOffsetsSizesAndStrides(st, sti, isEqualConstantIntOrValue)) return false; diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp index 06168d5ef2c7f..a8deeaf8a9eeb 100644 --- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp +++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp @@ -407,7 +407,7 @@ class SparseTensorToDimSizeConverter matchAndRewrite(tensor::DimOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { // Only rewrite annotated DimOp with constant index. - auto enc = getSparseTensorEncoding(op.source().getType()); + auto enc = getSparseTensorEncoding(op.getSource().getType()); if (!enc) return failure(); Optional index = op.getConstantIndex(); @@ -429,7 +429,7 @@ class SparseCastConverter : public OpConversionPattern { ConversionPatternRewriter &rewriter) const override { // Only rewrite identically annotated source/dest. auto encDst = getSparseTensorEncoding(op.getType()); - auto encSrc = getSparseTensorEncoding(op.source().getType()); + auto encSrc = getSparseTensorEncoding(op.getSource().getType()); if (!encDst || encDst != encSrc) return failure(); rewriter.replaceOp(op, adaptor.getOperands()); @@ -511,7 +511,7 @@ class SparseTensorConvertConverter : public OpConversionPattern { ConversionPatternRewriter &rewriter) const override { Location loc = op->getLoc(); Type resType = op.getType(); - Type srcType = op.source().getType(); + Type srcType = op.getSource().getType(); auto encDst = getSparseTensorEncoding(resType); auto encSrc = getSparseTensorEncoding(srcType); Value src = adaptor.getOperands()[0]; @@ -771,7 +771,7 @@ class SparseTensorLoadConverter : public OpConversionPattern { LogicalResult matchAndRewrite(LoadOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - if (op.hasInserts()) { + if (op.getHasInserts()) { // Finalize any pending insertions. StringRef name = "endInsert"; TypeRange noTp; @@ -790,7 +790,7 @@ class SparseTensorLexInsertConverter : public OpConversionPattern { LogicalResult matchAndRewrite(LexInsertOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - Type elemTp = op.tensor().getType().cast().getElementType(); + Type elemTp = op.getTensor().getType().cast().getElementType(); SmallString<12> name{"lexInsert", primaryTypeFunctionSuffix(elemTp)}; TypeRange noTp; replaceOpWithFuncCall(rewriter, op, name, noTp, adaptor.getOperands(), @@ -806,12 +806,12 @@ class SparseTensorExpandConverter : public OpConversionPattern { matchAndRewrite(ExpandOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { Location loc = op->getLoc(); - ShapedType srcType = op.tensor().getType().cast(); + ShapedType srcType = op.getTensor().getType().cast(); Type eltType = srcType.getElementType(); Type boolType = rewriter.getIntegerType(1); Type idxType = rewriter.getIndexType(); // All initialization should be done on entry of the loop nest. - rewriter.setInsertionPointAfter(op.tensor().getDefiningOp()); + rewriter.setInsertionPointAfter(op.getTensor().getDefiningOp()); // Determine the size for access expansion. auto enc = getSparseTensorEncoding(srcType); Value src = adaptor.getOperands()[0]; @@ -852,7 +852,7 @@ class SparseTensorCompressConverter : public OpConversionPattern { // all-zero/false by only iterating over the set elements, so the // complexity remains proportional to the sparsity of the expanded // access pattern. - Type elemTp = op.tensor().getType().cast().getElementType(); + Type elemTp = op.getTensor().getType().cast().getElementType(); SmallString<12> name{"expInsert", primaryTypeFunctionSuffix(elemTp)}; TypeRange noTp; replaceOpWithFuncCall(rewriter, op, name, noTp, adaptor.getOperands(), @@ -880,7 +880,7 @@ class SparseTensorOutConverter : public OpConversionPattern { matchAndRewrite(OutOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { Location loc = op->getLoc(); - ShapedType srcType = op.tensor().getType().cast(); + ShapedType srcType = op.getTensor().getType().cast(); // Convert to default permuted COO. Value src = adaptor.getOperands()[0]; auto encSrc = getSparseTensorEncoding(srcType); diff --git a/mlir/lib/Dialect/Tensor/Transforms/BufferizableOpInterfaceImpl.cpp b/mlir/lib/Dialect/Tensor/Transforms/BufferizableOpInterfaceImpl.cpp index 0e1617374ae31..f3056b9fd52bd 100644 --- a/mlir/lib/Dialect/Tensor/Transforms/BufferizableOpInterfaceImpl.cpp +++ b/mlir/lib/Dialect/Tensor/Transforms/BufferizableOpInterfaceImpl.cpp @@ -52,7 +52,7 @@ struct CastOpInterface auto castOp = cast(op); // The result buffer still has the old (pre-cast) type. - Value resultBuffer = getBuffer(rewriter, castOp.source(), options); + Value resultBuffer = getBuffer(rewriter, castOp.getSource(), options); auto sourceMemRefType = resultBuffer.getType().cast(); Attribute memorySpace = sourceMemRefType.getMemorySpace(); TensorType resultTensorType = @@ -108,7 +108,7 @@ struct CollapseShapeOpInterface const BufferizationOptions &options) const { auto collapseShapeOp = cast(op); RankedTensorType tensorResultType = collapseShapeOp.getResultType(); - Value buffer = getBuffer(rewriter, collapseShapeOp.src(), options); + Value buffer = getBuffer(rewriter, collapseShapeOp.getSrc(), options); auto bufferType = buffer.getType().cast(); if (tensorResultType.getRank() == 0) { @@ -135,7 +135,7 @@ struct CollapseShapeOpInterface } replaceOpWithNewBufferizedOp( - rewriter, op, resultType, buffer, collapseShapeOp.reassociation()); + rewriter, op, resultType, buffer, collapseShapeOp.getReassociation()); return success(); } @@ -148,8 +148,8 @@ struct CollapseShapeOpInterface // TODO: Create alloc_tensor ops during TensorCopyInsertion. AnalysisState analysisState(options); Value tensorAlloc = allocateTensorForShapedValue( - rewriter, op->getLoc(), collapseShapeOp.src(), - analysisState.isTensorYielded(collapseShapeOp.result())); + rewriter, op->getLoc(), collapseShapeOp.getSrc(), + analysisState.isTensorYielded(collapseShapeOp.getResult())); auto memrefType = MemRefType::get(collapseShapeOp.getSrcType().getShape(), collapseShapeOp.getSrcType().getElementType(), @@ -187,8 +187,9 @@ struct DimOpInterface LogicalResult bufferize(Operation *op, RewriterBase &rewriter, const BufferizationOptions &options) const { auto dimOp = cast(op); - auto v = getBuffer(rewriter, dimOp.source(), options); - replaceOpWithNewBufferizedOp(rewriter, op, v, dimOp.index()); + auto v = getBuffer(rewriter, dimOp.getSource(), options); + replaceOpWithNewBufferizedOp(rewriter, op, v, + dimOp.getIndex()); return success(); } }; @@ -223,7 +224,7 @@ struct ExpandShapeOpInterface const BufferizationOptions &options) const { auto expandShapeOp = cast(op); auto tensorResultType = expandShapeOp.getResultType(); - auto buffer = getBuffer(rewriter, expandShapeOp.src(), options); + auto buffer = getBuffer(rewriter, expandShapeOp.getSrc(), options); // Memref result type is inferred by the builder based on reassociation // indices and result shape. @@ -267,10 +268,10 @@ struct ExtractSliceOpInterface // Even if this op was decided to bufferize out-of-place, do not insert the // buffer copy yet. This is done later in this function. - auto srcMemref = getBuffer(rewriter, extractSliceOp.source(), options); + auto srcMemref = getBuffer(rewriter, extractSliceOp.getSource(), options); auto srcMemrefType = srcMemref.getType().cast(); auto dstTensorType = - extractSliceOp.result().getType().cast(); + extractSliceOp.getResult().getType().cast(); // Expand offsets, sizes and strides to the full rank to handle the // rank-reducing case. @@ -321,9 +322,9 @@ struct ExtractOpInterface LogicalResult bufferize(Operation *op, RewriterBase &rewriter, const BufferizationOptions &options) const { auto extractOp = cast(op); - Value srcMemref = getBuffer(rewriter, extractOp.tensor(), options); + Value srcMemref = getBuffer(rewriter, extractOp.getTensor(), options); replaceOpWithNewBufferizedOp(rewriter, op, srcMemref, - extractOp.indices()); + extractOp.getIndices()); return success(); } }; @@ -365,8 +366,8 @@ struct FromElementsOpInterface // TODO: Create alloc_tensor ops during TensorCopyInsertion. AnalysisState analysisState(options); Value tensorAlloc = allocateTensorForShapedValue( - rewriter, loc, fromElementsOp.result(), - analysisState.isTensorYielded(fromElementsOp.result()), + rewriter, loc, fromElementsOp.getResult(), + analysisState.isTensorYielded(fromElementsOp.getResult()), /*copy=*/false); auto memrefType = MemRefType::get(tensorType.getShape(), tensorType.getElementType()); @@ -374,15 +375,15 @@ struct FromElementsOpInterface op->getLoc(), memrefType, tensorAlloc); // Case: tensor<0xelem_type>. - if (fromElementsOp.elements().empty()) { + if (fromElementsOp.getElements().empty()) { replaceOpWithBufferizedValues(rewriter, op, buffer); return success(); } // Case: tensor. if (shape.empty()) { - rewriter.create(loc, fromElementsOp.elements().front(), - buffer); + rewriter.create( + loc, fromElementsOp.getElements().front(), buffer); replaceOpWithBufferizedValues(rewriter, op, buffer); return success(); } @@ -395,7 +396,7 @@ struct FromElementsOpInterface constants.push_back(rewriter.create(loc, i)); // Traverse all `elements` and create `memref.store` ops. - auto elementIt = fromElementsOp.elements().begin(); + auto elementIt = fromElementsOp.getElements().begin(); SmallVector indices(tensorType.getRank(), constants[0]); createStores(rewriter, loc, /*dim=*/0, buffer, shape, constants, elementIt, indices); @@ -418,8 +419,8 @@ struct GenerateOpInterface // TODO: Create alloc_tensor ops during TensorCopyInsertion. AnalysisState analysisState(options); Value tensorAlloc = allocateTensorForShapedValue( - rewriter, loc, generateOp.result(), - analysisState.isTensorYielded(generateOp.result()), + rewriter, loc, generateOp.getResult(), + analysisState.isTensorYielded(generateOp.getResult()), /*copy=*/false); auto memrefType = MemRefType::get(tensorType.getShape(), tensorType.getElementType()); @@ -435,10 +436,11 @@ struct GenerateOpInterface SmallVector upperBounds; int nextDynamicIndex = 0; for (int i = 0; i < rank; i++) { - Value upperBound = memrefType.isDynamicDim(i) - ? generateOp.dynamicExtents()[nextDynamicIndex++] - : rewriter.create( - loc, memrefType.getDimSize(i)); + Value upperBound = + memrefType.isDynamicDim(i) + ? generateOp.getDynamicExtents()[nextDynamicIndex++] + : rewriter.create( + loc, memrefType.getDimSize(i)); upperBounds.push_back(upperBound); } @@ -495,9 +497,9 @@ struct InsertOpInterface LogicalResult bufferize(Operation *op, RewriterBase &rewriter, const BufferizationOptions &options) const { auto insertOp = cast(op); - Value destMemref = getBuffer(rewriter, insertOp.dest(), options); - rewriter.create(insertOp.getLoc(), insertOp.scalar(), - destMemref, insertOp.indices()); + Value destMemref = getBuffer(rewriter, insertOp.getDest(), options); + rewriter.create(insertOp.getLoc(), insertOp.getScalar(), + destMemref, insertOp.getIndices()); replaceOpWithBufferizedValues(rewriter, op, destMemref); return success(); } @@ -519,7 +521,7 @@ static bool areEquivalentExtractSliceOps(const AnalysisState &state, if (!st || !sti) return false; if (sti != sti && - !state.areEquivalentBufferizedValues(st.source(), sti.dest())) + !state.areEquivalentBufferizedValues(st.getSource(), sti.getDest())) return false; if (!sameOffsetsSizesAndStrides(st, sti, isEqualConstantIntOrValue)) return false; @@ -636,8 +638,8 @@ struct InsertSliceOpInterface // is no memory write here.) if (uConflictingWrite == &insertSliceOp->getOpOperand(1) /*dest*/ && state.areEquivalentBufferizedValues(uRead->get(), - insertSliceOp.source()) && - hasMatchingExtractSliceOp(state, insertSliceOp.source(), + insertSliceOp.getSource()) && + hasMatchingExtractSliceOp(state, insertSliceOp.getSource(), insertSliceOp)) return true; @@ -653,7 +655,7 @@ struct InsertSliceOpInterface // TODO: be very loud about it or even consider failing the pass. auto insertSliceOp = cast(op); Location loc = insertSliceOp.getLoc(); - Value dstMemref = getBuffer(rewriter, insertSliceOp.dest(), options); + Value dstMemref = getBuffer(rewriter, insertSliceOp.getDest(), options); // Expand offsets, sizes and strides to the full rank to handle the // rank-reducing case. @@ -681,7 +683,7 @@ struct InsertSliceOpInterface // Copy tensor. If this tensor.insert_slice has a matching // tensor.extract_slice, the copy operation will eventually fold away. - auto srcMemref = getBuffer(rewriter, insertSliceOp.source(), options); + auto srcMemref = getBuffer(rewriter, insertSliceOp.getSource(), options); if (failed(options.createMemCpy(rewriter, loc, srcMemref, subView))) return failure(); @@ -712,7 +714,7 @@ struct RankOpInterface LogicalResult bufferize(Operation *op, RewriterBase &rewriter, const BufferizationOptions &options) const { auto rankOp = cast(op); - auto v = getBuffer(rewriter, rankOp.tensor(), options); + auto v = getBuffer(rewriter, rankOp.getTensor(), options); replaceOpWithNewBufferizedOp(rewriter, op, rankOp.getType(), v); return success(); @@ -748,8 +750,8 @@ struct ReshapeOpInterface LogicalResult bufferize(Operation *op, RewriterBase &rewriter, const BufferizationOptions &options) const { auto reshapeOp = cast(op); - Value srcBuffer = getBuffer(rewriter, reshapeOp.source(), options); - Value shapeBuffer = getBuffer(rewriter, reshapeOp.shape(), options); + Value srcBuffer = getBuffer(rewriter, reshapeOp.getSource(), options); + Value shapeBuffer = getBuffer(rewriter, reshapeOp.getShape(), options); auto resultTensorType = reshapeOp.getResult().getType().cast(); auto resultMemRefType = getMemRefType(resultTensorType, options); replaceOpWithNewBufferizedOp( diff --git a/mlir/lib/Dialect/X86Vector/IR/X86VectorDialect.cpp b/mlir/lib/Dialect/X86Vector/IR/X86VectorDialect.cpp index 7b70e53a6e9c3..ac21f1714689d 100644 --- a/mlir/lib/Dialect/X86Vector/IR/X86VectorDialect.cpp +++ b/mlir/lib/Dialect/X86Vector/IR/X86VectorDialect.cpp @@ -29,13 +29,13 @@ void x86vector::X86VectorDialect::initialize() { } LogicalResult x86vector::MaskCompressOp::verify() { - if (src() && constant_src()) + if (getSrc() && getConstantSrc()) return emitError("cannot use both src and constant_src"); - if (src() && (src().getType() != dst().getType())) + if (getSrc() && (getSrc().getType() != getDst().getType())) return emitError("failed to verify that src and dst have same type"); - if (constant_src() && (constant_src()->getType() != dst().getType())) + if (getConstantSrc() && (getConstantSrc()->getType() != getDst().getType())) return emitError( "failed to verify that constant_src and dst have same type"); diff --git a/mlir/lib/Dialect/X86Vector/Transforms/LegalizeForLLVMExport.cpp b/mlir/lib/Dialect/X86Vector/Transforms/LegalizeForLLVMExport.cpp index bdd8d1fd31e60..4df05b6b1d0f9 100644 --- a/mlir/lib/Dialect/X86Vector/Transforms/LegalizeForLLVMExport.cpp +++ b/mlir/lib/Dialect/X86Vector/Transforms/LegalizeForLLVMExport.cpp @@ -22,11 +22,11 @@ using namespace mlir::x86vector; /// Extracts the "main" vector element type from the given X86Vector operation. template static Type getSrcVectorElementType(OpTy op) { - return op.src().getType().template cast().getElementType(); + return op.getSrc().getType().template cast().getElementType(); } template <> Type getSrcVectorElementType(Vp2IntersectOp op) { - return op.a().getType().template cast().getElementType(); + return op.getA().getType().template cast().getElementType(); } namespace { @@ -70,21 +70,21 @@ struct MaskCompressOpConversion LogicalResult matchAndRewrite(MaskCompressOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - auto opType = adaptor.a().getType(); + auto opType = adaptor.getA().getType(); Value src; - if (op.src()) { - src = adaptor.src(); - } else if (op.constant_src()) { + if (op.getSrc()) { + src = adaptor.getSrc(); + } else if (op.getConstantSrc()) { src = rewriter.create(op.getLoc(), opType, - op.constant_srcAttr()); + op.getConstantSrcAttr()); } else { Attribute zeroAttr = rewriter.getZeroAttr(opType); src = rewriter.create(op->getLoc(), opType, zeroAttr); } - rewriter.replaceOpWithNewOp(op, opType, adaptor.a(), - src, adaptor.k()); + rewriter.replaceOpWithNewOp(op, opType, adaptor.getA(), + src, adaptor.getK()); return success(); } @@ -96,8 +96,8 @@ struct RsqrtOpConversion : public ConvertOpToLLVMPattern { LogicalResult matchAndRewrite(RsqrtOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - auto opType = adaptor.a().getType(); - rewriter.replaceOpWithNewOp(op, opType, adaptor.a()); + auto opType = adaptor.getA().getType(); + rewriter.replaceOpWithNewOp(op, opType, adaptor.getA()); return success(); } }; @@ -108,14 +108,14 @@ struct DotOpConversion : public ConvertOpToLLVMPattern { LogicalResult matchAndRewrite(DotOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { - auto opType = adaptor.a().getType(); + auto opType = adaptor.getA().getType(); Type llvmIntType = IntegerType::get(&getTypeConverter()->getContext(), 8); // Dot product of all elements, broadcasted to all elements. auto attr = rewriter.getI8IntegerAttr(static_cast(0xff)); Value scale = rewriter.create(op.getLoc(), llvmIntType, attr); - rewriter.replaceOpWithNewOp(op, opType, adaptor.a(), adaptor.b(), - scale); + rewriter.replaceOpWithNewOp(op, opType, adaptor.getA(), + adaptor.getB(), scale); return success(); } }; diff --git a/mlir/lib/Target/Cpp/TranslateToCpp.cpp b/mlir/lib/Target/Cpp/TranslateToCpp.cpp index 4eaa686e56d8f..f037b7d8febf7 100644 --- a/mlir/lib/Target/Cpp/TranslateToCpp.cpp +++ b/mlir/lib/Target/Cpp/TranslateToCpp.cpp @@ -217,7 +217,7 @@ static LogicalResult printConstantOp(CppEmitter &emitter, Operation *operation, static LogicalResult printOperation(CppEmitter &emitter, emitc::ConstantOp constantOp) { Operation *operation = constantOp.getOperation(); - Attribute value = constantOp.value(); + Attribute value = constantOp.getValue(); return printConstantOp(emitter, operation, value); } @@ -225,7 +225,7 @@ static LogicalResult printOperation(CppEmitter &emitter, static LogicalResult printOperation(CppEmitter &emitter, emitc::VariableOp variableOp) { Operation *operation = variableOp.getOperation(); - Attribute value = variableOp.value(); + Attribute value = variableOp.getValue(); return printConstantOp(emitter, operation, value); } @@ -330,7 +330,7 @@ static LogicalResult printOperation(CppEmitter &emitter, emitc::CallOp callOp) { if (failed(emitter.emitAssignPrefix(op))) return failure(); - os << callOp.callee(); + os << callOp.getCallee(); auto emitArgs = [&](Attribute attr) -> LogicalResult { if (auto t = attr.dyn_cast()) { @@ -352,9 +352,10 @@ static LogicalResult printOperation(CppEmitter &emitter, emitc::CallOp callOp) { return success(); }; - if (callOp.template_args()) { + if (callOp.getTemplateArgs()) { os << "<"; - if (failed(interleaveCommaWithError(*callOp.template_args(), os, emitArgs))) + if (failed( + interleaveCommaWithError(*callOp.getTemplateArgs(), os, emitArgs))) return failure(); os << ">"; } @@ -362,8 +363,9 @@ static LogicalResult printOperation(CppEmitter &emitter, emitc::CallOp callOp) { os << "("; LogicalResult emittedArgs = - callOp.args() ? interleaveCommaWithError(*callOp.args(), os, emitArgs) - : emitter.emitOperands(op); + callOp.getArgs() + ? interleaveCommaWithError(*callOp.getArgs(), os, emitArgs) + : emitter.emitOperands(op); if (failed(emittedArgs)) return failure(); os << ")"; @@ -377,7 +379,7 @@ static LogicalResult printOperation(CppEmitter &emitter, if (failed(emitter.emitAssignPrefix(op))) return failure(); - os << applyOp.applicableOperator(); + os << applyOp.getApplicableOperator(); os << emitter.getOrCreateName(applyOp.getOperand()); return success(); @@ -403,10 +405,10 @@ static LogicalResult printOperation(CppEmitter &emitter, raw_ostream &os = emitter.ostream(); os << "#include "; - if (includeOp.is_standard_include()) - os << "<" << includeOp.include() << ">"; + if (includeOp.getIsStandardInclude()) + os << "<" << includeOp.getInclude() << ">"; else - os << "\"" << includeOp.include() << "\""; + os << "\"" << includeOp.getInclude() << "\""; return success(); }