[llvm-branch-commits] [flang] [mlir] [mlir] Migrate aggregate builders to explicit properties (PR #219195)
Mehdi Amini via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Aug 27 07:16:25 PDT 2026
https://github.com/joker-eph updated https://github.com/llvm/llvm-project/pull/219195
>From 545dcb648551098e0922f19d80fd8e54dafb7084 Mon Sep 17 00:00:00 2001
From: Mehdi Amini <joker.eph at gmail.com>
Date: Thu, 20 Aug 2026 07:49:23 -0700
Subject: [PATCH] [mlir] Migrate aggregate builders to explicit properties
Pass typed property structs and discardable attributes separately at ODS
aggregate builder call sites.
Assisted-by: Codex
---
flang/lib/Lower/OpenACC.cpp | 9 +-
.../Optimizer/Builder/CUDAIntrinsicCall.cpp | 3 +-
flang/lib/Optimizer/Builder/IntrinsicCall.cpp | 7 +-
.../Optimizer/Builder/PPCIntrinsicCall.cpp | 13 +-
flang/lib/Optimizer/CodeGen/CodeGen.cpp | 108 ++++++++++++----
flang/lib/Optimizer/CodeGen/TargetRewrite.cpp | 3 +-
mlir/docs/DefiningDialects/Operations.md | 18 +--
.../ArithCommon/AttrToLLVMConverter.h | 27 +++-
.../mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td | 11 +-
.../include/mlir/Dialect/LLVMIR/LLVMOpBase.td | 11 +-
.../mlir/Dialect/MemRef/IR/MemRefOps.td | 8 +-
.../mlir/Dialect/Tensor/IR/TensorOps.td | 8 +-
mlir/include/mlir/TableGen/Operator.h | 15 +++
.../AMDGPUToROCDL/AMDGPUToROCDL.cpp | 83 ++++++++-----
.../ArithToAMDGPU/ArithToAMDGPU.cpp | 20 ++-
.../Conversion/ArithToEmitC/ArithToEmitC.cpp | 5 +-
mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp | 12 +-
.../Conversion/MathToFuncs/MathToFuncs.cpp | 5 +-
mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp | 65 ++++++----
mlir/lib/Conversion/MathToLibm/MathToLibm.cpp | 13 +-
mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp | 4 +-
.../Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp | 2 +-
.../Conversion/RaiseWasm/RaiseWasmMLIR.cpp | 27 +++-
mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp | 3 +-
.../Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp | 94 ++++++++++----
.../Conversion/TosaToLinalg/TosaToLinalg.cpp | 106 ++++++++++------
.../TosaToLinalg/TosaToLinalgNamed.cpp | 4 +-
.../Conversion/VectorToGPU/VectorToGPU.cpp | 4 +-
.../Conversion/VectorToSCF/VectorToSCF.cpp | 6 +-
.../VectorToXeGPU/VectorToXeGPU.cpp | 6 +-
.../Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp | 4 +-
.../AMDGPU/Transforms/EmulateAtomics.cpp | 40 ++++--
mlir/lib/Dialect/Arith/IR/ArithOps.cpp | 7 +-
.../Transforms/EmulateUnsupportedFloats.cpp | 3 +-
.../Dialect/Arith/Transforms/ExpandOps.cpp | 3 +-
.../Transforms/UnsignedWhenEquivalent.cpp | 9 +-
mlir/lib/Dialect/Arith/Utils/Utils.cpp | 11 +-
.../GPU/Transforms/DecomposeMemRefs.cpp | 2 +-
.../Linalg/Transforms/Vectorization.cpp | 3 +-
.../Transforms/ExtendToSupportedTypes.cpp | 4 +-
.../Transforms/PolynomialApproximation.cpp | 9 +-
mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp | 14 ++-
.../Dialect/MemRef/Transforms/MultiBuffer.cpp | 2 +
mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp | 2 +-
.../Dialect/OpenACC/Transforms/ACCCGToGPU.cpp | 14 ++-
mlir/lib/Dialect/Shape/IR/Shape.cpp | 43 ++++---
.../Transforms/SparseVectorization.cpp | 20 +--
.../lib/Dialect/SparseTensor/Utils/Merger.cpp | 3 +-
mlir/lib/Dialect/Tensor/IR/TensorOps.cpp | 8 +-
.../Tosa/Transforms/TosaNarrowTypes.cpp | 3 +-
.../Vector/Transforms/LowerVectorContract.cpp | 3 +-
.../XeGPU/Transforms/XeGPUPropagateLayout.cpp | 5 +-
.../Transforms/XeGPUSgToLaneDistribute.cpp | 5 +-
.../Dialect/XeGPU/Transforms/XeGPUUnroll.cpp | 25 +++-
.../Transforms/XeGPUWgToSgDistribute.cpp | 4 +-
mlir/lib/TableGen/Operator.cpp | 29 +++++
.../SPIRV/Deserialization/DeserializeOps.cpp | 13 +-
mlir/lib/Tools/PDLL/CodeGen/MLIRGen.cpp | 3 +-
mlir/test/Dialect/SPIRV/IR/target-env.mlir | 3 +-
.../Dialect/Tosa/tosa-narrow-i64-to-i32.mlir | 4 +-
.../lib/Dialect/SPIRV/TestAvailability.cpp | 21 +++-
mlir/test/lib/Dialect/Test/TestPatterns.cpp | 3 +-
mlir/test/mlir-tblgen/op-decl-and-defs.td | 28 ++++-
mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp | 75 +++++-------
mlir/tools/mlir-tblgen/SPIRVUtilsGen.cpp | 40 +++++-
.../tools/tblgen-to-irdl/OpDefinitionsGen.cpp | 5 +-
.../Dialect/OpenACC/OpenACCOpsTest.cpp | 30 +++--
.../Dialect/OpenACC/OpenACCUtilsLoopTest.cpp | 14 ++-
.../Dialect/OpenACC/OpenACCUtilsTest.cpp | 115 +++++++-----------
mlir/unittests/TableGen/OpBuildGen.cpp | 18 ++-
70 files changed, 914 insertions(+), 433 deletions(-)
diff --git a/flang/lib/Lower/OpenACC.cpp b/flang/lib/Lower/OpenACC.cpp
index 30364d0ae9231..824ac1d0d24a7 100644
--- a/flang/lib/Lower/OpenACC.cpp
+++ b/flang/lib/Lower/OpenACC.cpp
@@ -150,7 +150,8 @@ createDataEntryOp(fir::FirOpBuilder &builder, mlir::Location loc,
addOperands(operands, operandSegments, bounds);
addOperands(operands, operandSegments, async);
- Op op = Op::create(builder, loc, retTy, operands);
+ Op op = Op::create(builder, loc, mlir::TypeRange{retTy}, operands,
+ typename Op::Properties{});
op.setNameAttr(builder.getStringAttr(name.str()));
op.setStructured(structured);
op.setImplicit(implicit);
@@ -218,8 +219,8 @@ static Op
createSimpleOp(fir::FirOpBuilder &builder, mlir::Location loc,
const llvm::SmallVectorImpl<mlir::Value> &operands,
const llvm::SmallVectorImpl<int32_t> &operandSegments) {
- llvm::ArrayRef<mlir::Type> argTy;
- Op op = Op::create(builder, loc, argTy, operands);
+ Op op = Op::create(builder, loc, mlir::TypeRange{}, operands,
+ typename Op::Properties{});
op->setAttr(Op::getOperandSegmentSizeAttr(),
builder.getDenseI32ArrayAttr(operandSegments));
return op;
@@ -1301,7 +1302,7 @@ createRegionOp(fir::FirOpBuilder &builder, mlir::Location loc,
llvm::SmallVector<mlir::Type> retTy = {},
mlir::Value yieldValue = {}, mlir::TypeRange argsTy = {},
llvm::SmallVector<mlir::Location> locs = {}) {
- Op op = Op::create(builder, loc, retTy, operands);
+ Op op = Op::create(builder, loc, retTy, operands, typename Op::Properties{});
builder.createBlock(&op.getRegion(), op.getRegion().end(), argsTy, locs);
mlir::Block &block = op.getRegion().back();
builder.setInsertionPointToStart(&block);
diff --git a/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp b/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
index ca200ac2cd02a..c6866f6d2f136 100644
--- a/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
@@ -977,7 +977,8 @@ CUDAIntrinsicLibrary::genBarrierArrive(mlir::Type resultType,
assert(args.size() == 1);
mlir::Value barrier = convertPtrToNVVMSpace(
builder, loc, args[0], mlir::NVVM::NVVMMemorySpace::Shared);
- return mlir::NVVM::MBarrierArriveOp::create(builder, loc, resultType, barrier)
+ return mlir::NVVM::MBarrierArriveOp::create(builder, loc, resultType, barrier,
+ /*count=*/nullptr)
.getResult(0);
}
diff --git a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
index 726319b11a720..9f9832f790827 100644
--- a/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/IntrinsicCall.cpp
@@ -1113,7 +1113,7 @@ mlir::Value genMathOp(fir::FirOpBuilder &builder, mlir::Location loc,
LLVM_DEBUG(llvm::dbgs() << "Generating '" << mathLibFuncName
<< "' operation with type ";
mathLibFuncType.dump(); llvm::dbgs() << "\n");
- result = T::create(builder, loc, args);
+ result = T::create(builder, loc, args, typename T::Properties{});
}
LLVM_DEBUG(result.dump(); llvm::dbgs() << "\n");
return result;
@@ -1151,12 +1151,13 @@ mlir::Value genComplexMathOp(fir::FirOpBuilder &builder, mlir::Location loc,
// the argument types for an operation
if constexpr (T::template hasTrait<
mlir::OpTrait::SameOperandsAndResultType>()) {
- result = T::create(builder, loc, args);
+ result = T::create(builder, loc, args, typename T::Properties{});
result = builder.createConvert(loc, mathLibFuncType.getResult(0), result);
} else {
auto complexTy = mlir::cast<mlir::ComplexType>(mathLibFuncType.getInput(0));
auto realTy = complexTy.getElementType();
- result = T::create(builder, loc, realTy, args);
+ result = T::create(builder, loc, mlir::TypeRange{realTy}, args,
+ typename T::Properties{});
result = builder.createConvert(loc, mathLibFuncType.getResult(0), result);
}
diff --git a/flang/lib/Optimizer/Builder/PPCIntrinsicCall.cpp b/flang/lib/Optimizer/Builder/PPCIntrinsicCall.cpp
index cf8c7ce4f8d4e..2e52939576538 100644
--- a/flang/lib/Optimizer/Builder/PPCIntrinsicCall.cpp
+++ b/flang/lib/Optimizer/Builder/PPCIntrinsicCall.cpp
@@ -1877,8 +1877,10 @@ fir::ExtendedValue PPCIntrinsicLibrary::genVecLdNoCallGrp(
const auto triple{fir::getTargetTriple(builder.getModule())};
// Need to get align 1.
- auto result{fir::LoadOp::create(builder, loc, mlirTy, addr,
- getAlignmentAttr(builder, 1))};
+ auto result{fir::LoadOp::create(
+ builder, loc, mlir::TypeRange{mlirTy}, mlir::ValueRange{addr},
+ fir::LoadOp::Properties{},
+ llvm::ArrayRef<mlir::NamedAttribute>{getAlignmentAttr(builder, 1)})};
if ((vop == VecOp::Xl && isBEVecElemOrderOnLE()) ||
(vop == VecOp::Xlbe && triple.isLittleEndian()))
return builder.createConvert(
@@ -2999,9 +3001,10 @@ void PPCIntrinsicLibrary::genVecXStore(
default:
assert(false && "Invalid vector operation for generator");
}
- fir::StoreOp::create(builder, loc, mlir::TypeRange{},
- mlir::ValueRange{src, trg},
- getAlignmentAttr(builder, 1));
+ fir::StoreOp::create(
+ builder, loc, mlir::TypeRange{}, mlir::ValueRange{src, trg},
+ fir::StoreOp::Properties{},
+ llvm::ArrayRef<mlir::NamedAttribute>{getAlignmentAttr(builder, 1)});
}
} // namespace fir
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index c5360d468bdd7..a2e7ce7403b6e 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -157,6 +157,44 @@ addLLVMOpBundleAttrs(mlir::ConversionPatternRewriter &rewriter,
return newAttrs;
}
+template <typename Op>
+struct BuilderAttributes {
+ typename Op::Properties properties{};
+ llvm::SmallVector<mlir::NamedAttribute> discardableAttributes;
+};
+
+template <typename Op>
+static BuilderAttributes<Op>
+splitBuilderAttributes(mlir::ConversionPatternRewriter &rewriter,
+ llvm::ArrayRef<mlir::NamedAttribute> attributes) {
+ BuilderAttributes<Op> result;
+ Op::populateDefaultProperties(
+ mlir::OperationName(Op::getOperationName(), rewriter.getContext()),
+ result.properties);
+ mlir::LogicalResult converted = Op::setPropertiesFromAttr(
+ result.properties, rewriter.getDictionaryAttr(attributes), [&]() {
+ return mlir::emitError(rewriter.getUnknownLoc(),
+ "failed to convert operation properties");
+ });
+ assert(mlir::succeeded(converted) && "failed to convert properties");
+ (void)converted;
+
+ for (mlir::NamedAttribute attr : attributes) {
+ if (!Op::getInherentAttr(rewriter.getContext(), result.properties,
+ attr.getName().getValue()))
+ result.discardableAttributes.push_back(attr);
+ }
+ return result;
+}
+
+static BuilderAttributes<mlir::LLVM::CallOp>
+getLLVMCallBuilderAttributes(mlir::ConversionPatternRewriter &rewriter,
+ llvm::ArrayRef<mlir::NamedAttribute> attributes,
+ int32_t numCallOperands) {
+ return splitBuilderAttributes<mlir::LLVM::CallOp>(
+ rewriter, addLLVMOpBundleAttrs(rewriter, attributes, numCallOperands));
+}
+
namespace {
// Replaces an existing operation with an AddressOfOp or an AddrSpaceCastOp
@@ -347,9 +385,11 @@ struct AllocaOpConversion : public fir::FIROpConversion<fir::AllocaOp> {
emitError(loc, "did not find allocation function");
mlir::NamedAttribute attr = rewriter.getNamedAttr(
"callee", mlir::SymbolRefAttr::get(memSizeFn));
+ auto builderAttrs =
+ getLLVMCallBuilderAttributes(rewriter, {attr}, lenParams.size());
auto call = mlir::LLVM::CallOp::create(
- rewriter, loc, ity, lenParams,
- addLLVMOpBundleAttrs(rewriter, {attr}, lenParams.size()));
+ rewriter, loc, mlir::TypeRange{ity}, lenParams,
+ builderAttrs.properties, builderAttrs.discardableAttributes);
size = call.getResult();
llvmObjectType = ::getI8Type(alloc.getContext());
} else {
@@ -720,10 +760,13 @@ struct CallOpConversion : public fir::FIROpConversion<fir::CallOp> {
// Convert arith::FastMathFlagsAttr to LLVM::FastMathFlagsAttr.
mlir::arith::AttrConvertFastMathToLLVM<fir::CallOp, mlir::LLVM::CallOp>
attrConvert(call);
+ auto builderAttrs = getLLVMCallBuilderAttributes(
+ rewriter, attrConvert.getAttrs(), adaptor.getOperands().size());
+ builderAttrs.properties.fastmathFlags =
+ attrConvert.getProperties().fastmathFlags;
auto llvmCall = rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>(
- call, resultTys, adaptor.getOperands(),
- addLLVMOpBundleAttrs(rewriter, attrConvert.getAttrs(),
- adaptor.getOperands().size()));
+ call, resultTys, adaptor.getOperands(), builderAttrs.properties,
+ builderAttrs.discardableAttributes);
if (mlir::ArrayAttr argAttrsArray = call.getArgAttrsAttr()) {
// sret and byval type needs to be converted.
auto convertTypeAttr = [&](const mlir::NamedAttribute &attr) {
@@ -1490,12 +1533,14 @@ struct AllocMemOpConversion : public fir::FIROpConversion<fir::AllocMemOp> {
mlir::LLVM::StoreOp::create(rewriter, loc, nullPtr, memptr);
heap->setAttr("callee", getPosixMemalign(heap, rewriter, mallocTy,
this->options));
- mlir::LLVM::CallOp::create(
- rewriter, loc,
- mlir::TypeRange{
- mlir::IntegerType::get(rewriter.getContext(), 32)},
- mlir::ValueRange{memptr, alignVal, size},
- addLLVMOpBundleAttrs(rewriter, heap->getAttrs(), 3));
+ auto builderAttrs =
+ getLLVMCallBuilderAttributes(rewriter, heap->getAttrs(), 3);
+ mlir::LLVM::CallOp::create(rewriter, loc,
+ mlir::TypeRange{mlir::IntegerType::get(
+ rewriter.getContext(), 32)},
+ mlir::ValueRange{memptr, alignVal, size},
+ builderAttrs.properties,
+ builderAttrs.discardableAttributes);
mlir::Value newPtr =
mlir::LLVM::LoadOp::create(rewriter, loc, ptrTy, memptr);
rewriter.replaceOp(heap, newPtr);
@@ -1513,18 +1558,23 @@ struct AllocMemOpConversion : public fir::FIROpConversion<fir::AllocMemOp> {
rewriter, loc, mallocTy, sizePlus, notAlignMinusOne);
heap->setAttr("callee",
getAlignedAlloc(heap, rewriter, mallocTy, this->options));
+ auto builderAttrs =
+ getLLVMCallBuilderAttributes(rewriter, heap->getAttrs(), 2);
rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>(
- heap, ::getLlvmPtrType(heap.getContext()),
- mlir::ValueRange{alignVal, roundedSize},
- addLLVMOpBundleAttrs(rewriter, heap->getAttrs(), 2));
+ heap, mlir::TypeRange{::getLlvmPtrType(heap.getContext())},
+ mlir::ValueRange{alignVal, roundedSize}, builderAttrs.properties,
+ builderAttrs.discardableAttributes);
return mlir::success();
}
}
heap->setAttr("callee", getMalloc(heap, rewriter, mallocTy, this->options));
+ auto builderAttrs =
+ getLLVMCallBuilderAttributes(rewriter, heap->getAttrs(), 1);
rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>(
- heap, ::getLlvmPtrType(heap.getContext()), size,
- addLLVMOpBundleAttrs(rewriter, heap->getAttrs(), 1));
+ heap, mlir::TypeRange{::getLlvmPtrType(heap.getContext())},
+ mlir::ValueRange{size}, builderAttrs.properties,
+ builderAttrs.discardableAttributes);
return mlir::success();
}
@@ -1594,10 +1644,12 @@ struct FreeMemOpConversion : public fir::FIROpConversion<fir::FreeMemOp> {
mlir::ConversionPatternRewriter &rewriter) const override {
mlir::Location loc = freemem.getLoc();
freemem->setAttr("callee", getFree(freemem, rewriter, this->options));
- mlir::LLVM::CallOp::create(
- rewriter, loc, mlir::TypeRange{},
- mlir::ValueRange{adaptor.getHeapref()},
- addLLVMOpBundleAttrs(rewriter, freemem->getAttrs(), 1));
+ auto builderAttrs =
+ getLLVMCallBuilderAttributes(rewriter, freemem->getAttrs(), 1);
+ mlir::LLVM::CallOp::create(rewriter, loc, mlir::TypeRange{},
+ mlir::ValueRange{adaptor.getHeapref()},
+ builderAttrs.properties,
+ builderAttrs.discardableAttributes);
rewriter.eraseOp(freemem);
return mlir::success();
}
@@ -3573,10 +3625,11 @@ struct FieldIndexOpConversion : public fir::FIROpConversion<fir::FieldIndexOp> {
mlir::NamedAttribute callAttr = rewriter.getNamedAttr("callee", symAttr);
mlir::NamedAttribute fieldAttr = rewriter.getNamedAttr(
"field", mlir::IntegerAttr::get(lowerTy().indexType(), index));
+ auto builderAttrs = getLLVMCallBuilderAttributes(
+ rewriter, {callAttr, fieldAttr}, adaptor.getOperands().size());
rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>(
- field, lowerTy().offsetType(), adaptor.getOperands(),
- addLLVMOpBundleAttrs(rewriter, {callAttr, fieldAttr},
- adaptor.getOperands().size()));
+ field, mlir::TypeRange{lowerTy().offsetType()}, adaptor.getOperands(),
+ builderAttrs.properties, builderAttrs.discardableAttributes);
return mlir::success();
}
@@ -3941,9 +3994,12 @@ struct LoadOpConversion : public fir::FIROpConversion<fir::LoadOp> {
rewriter.replaceOp(load, newBoxStorage);
} else {
- mlir::LLVM::LoadOp loadOp =
- mlir::LLVM::LoadOp::create(rewriter, load.getLoc(), llvmLoadTy,
- adaptor.getOperands(), load->getAttrs());
+ auto builderAttrs = splitBuilderAttributes<mlir::LLVM::LoadOp>(
+ rewriter, load->getAttrs());
+ mlir::LLVM::LoadOp loadOp = mlir::LLVM::LoadOp::create(
+ rewriter, load.getLoc(), mlir::TypeRange{llvmLoadTy},
+ adaptor.getOperands(), builderAttrs.properties,
+ builderAttrs.discardableAttributes);
loadOp.setVolatile_(isVolatile);
if (std::optional<mlir::ArrayAttr> optionalTag = load.getTbaa())
loadOp.setTBAATags(*optionalTag);
diff --git a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
index 8070877554ab1..f8eb8087a9b90 100644
--- a/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
+++ b/flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
@@ -567,7 +567,8 @@ class TargetRewrite : public fir::impl::TargetRewritePassBase<TargetRewrite> {
newOpers[0].setType(mlir::FunctionType::get(
callOp.getContext(),
mlir::TypeRange{newInTypes}.drop_front(dropFront), newResTys));
- newCall = fir::CallOp::create(*rewriter, loc, newResTys, newOpers);
+ newCall = fir::CallOp::create(*rewriter, loc, newResTys, newOpers,
+ fir::CallOp::Properties{});
}
newCall.setFastmathAttr(callOp.getFastmathAttr());
// Always set ABI argument attributes on call operations, even when
diff --git a/mlir/docs/DefiningDialects/Operations.md b/mlir/docs/DefiningDialects/Operations.md
index 7388d45ec9614..dc64edd75cd59 100644
--- a/mlir/docs/DefiningDialects/Operations.md
+++ b/mlir/docs/DefiningDialects/Operations.md
@@ -540,16 +540,20 @@ the same form regardless of the exact op. This is particularly useful for
implementing declarative pattern rewrites.
For operations with non-empty properties, the aggregate builder that takes a
-mixed `attributes` array partitions the array using the operation's statically
-known inherent-attribute and property names. It converts that subset into
-`Properties` and places only the remaining discardable attributes in
-`OperationState::attributes`. Defaults and result-type inference therefore
-observe the populated properties before the operation is created. Operations
-with empty properties retain the ordinary aggregate attribute builder.
+mixed `attributes` array is deprecated. Use the overload that takes a typed
+`Properties` structure and a separate `discardableAttributes` array instead.
+The deprecated overload remains available for compatibility: it partitions the
+mixed array using the operation's statically known inherent-attribute and
+property names, converts that subset into `Properties`, and places only the
+remaining discardable attributes in `OperationState::attributes`. Defaults and
+result-type inference therefore observe the populated properties before the
+operation is created. Operations with empty properties retain the ordinary
+aggregate attribute builder without a deprecation.
This applies to all aggregate builder variants, including builders with
explicit or inferred result types and builders that derive result types from
-operands or the first attribute.
+operands or the first attribute. The overload taking `Properties` and
+`discardableAttributes` is not deprecated.
The third and fourth forms are good for use in manually written code, given that
they provide better guarantee via signatures.
diff --git a/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h b/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h
index feb74c86e349f..a145fd64aff14 100644
--- a/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h
+++ b/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h
@@ -51,7 +51,8 @@ getLLVMDefaultFPExceptionBehavior(MLIRContext &context);
template <typename SourceOp, typename TargetOp>
class AttrConvertFastMathToLLVM {
public:
- AttrConvertFastMathToLLVM(SourceOp srcOp) {
+ AttrConvertFastMathToLLVM(SourceOp srcOp)
+ : context(srcOp.getOperation()->getContext()) {
// Copy the source attributes.
convertedAttr = NamedAttrList{srcOp->getAttrs()};
// Get the name of the arith fastmath attribute.
@@ -61,15 +62,33 @@ class AttrConvertFastMathToLLVM {
convertedAttr.erase(arithFMFAttrName));
if (arithFMFAttr) {
StringRef targetAttrName = TargetOp::getFastmathAttrName();
- convertedAttr.set(targetAttrName,
- convertArithFastMathAttrToLLVM(arithFMFAttr));
+ Builder builder(context);
+ propertiesAttr = builder.getDictionaryAttr(builder.getNamedAttr(
+ targetAttrName, convertArithFastMathAttrToLLVM(arithFMFAttr)));
}
}
ArrayRef<NamedAttribute> getAttrs() const { return convertedAttr.getAttrs(); }
- Attribute getPropAttr() const { return {}; }
+ Attribute getPropAttr() const { return propertiesAttr; }
+
+ typename TargetOp::Properties getProperties() const {
+ typename TargetOp::Properties properties{};
+ TargetOp::populateDefaultProperties(
+ OperationName(TargetOp::getOperationName(), context), properties);
+ if (propertiesAttr) {
+ LogicalResult result =
+ TargetOp::setPropertiesFromAttr(properties, propertiesAttr, [&]() {
+ return emitError(UnknownLoc::get(context));
+ });
+ assert(succeeded(result) && "failed to convert target properties");
+ (void)result;
+ }
+ return properties;
+ }
private:
+ MLIRContext *context;
NamedAttrList convertedAttr;
+ DictionaryAttr propertiesAttr;
};
// Attribute converter that populates a NamedAttrList by removing the overflow
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
index f40f3f5ebda93..83f56ba59306c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
@@ -501,8 +501,17 @@ class LLVM_ConstrainedIntr<string mnem, int numArgs,
$_builder.getNamedAttr($_qualCppClassName::getRoundingModeAttrName(),
roundingModeAttr));
}], true : "") # [{
+ $_qualCppClassName::Properties properties{};
+ $_qualCppClassName::populateDefaultProperties(
+ OperationName($_qualCppClassName::getOperationName(),
+ $_builder.getContext()), properties);
+ if (failed($_qualCppClassName::setPropertiesFromAttr(
+ properties, $_builder.getDictionaryAttr(mlirAttrs),
+ [&]() { return emitError($_location); }))) {
+ return failure();
+ }
auto op = $_qualCppClassName::create($_builder, $_location,
- $_resultType, mlirOperands, mlirAttrs);
+ $_resultType, mlirOperands, properties, {});
moduleImport.setFastmathFlagsAttr(inst, op);
$res = op;
}];
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
index a787840663171..da3036dc3187c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
@@ -385,8 +385,17 @@ class LLVM_IntrOpBase<Dialect dialect, string opName, string enumName,
}
SmallVector<Type> resultTypes =
}] # !if(!gt(numResults, 0), "{$_resultType};", "{};") # [{
+ $_qualCppClassName::Properties properties{};
+ $_qualCppClassName::populateDefaultProperties(
+ OperationName($_qualCppClassName::getOperationName(),
+ $_builder.getContext()), properties);
+ if (failed($_qualCppClassName::setPropertiesFromAttr(
+ properties, $_builder.getDictionaryAttr(mlirAttrs),
+ [&]() { return emitError($_location); }))) {
+ return failure();
+ }
auto op = $_qualCppClassName::create($_builder,
- $_location, resultTypes, mlirOperands, mlirAttrs);
+ $_location, resultTypes, mlirOperands, properties, {});
}];
string baseMlirBuilderArgAndResultAttrs = [{
moduleImport.convertArgAndResultAttrs(
diff --git a/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td b/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td
index 810eda17268ac..86e466c46c228 100644
--- a/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td
+++ b/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td
@@ -2027,9 +2027,11 @@ def MemRef_CollapseShapeOp : MemRef_ReassociativeReshapeOp<"collapse_shape", [
"ArrayRef<ReassociationIndices>":$reassociation,
CArg<"ArrayRef<NamedAttribute>", "{}">:$attrs),
[{
- $_state.addAttribute("reassociation",
- getReassociationIndicesAttribute($_builder, reassociation));
- build($_builder, $_state, resultType, src, attrs);
+ buildPropertiesAndDiscardableAttributes($_state, attrs);
+ $_state.getOrAddProperties<Properties>().reassociation =
+ getReassociationIndicesAttribute($_builder, reassociation);
+ $_state.addOperands(src);
+ $_state.addTypes(resultType);
}]>,
OpBuilder<(ins "Type":$resultType, "Value":$src,
"ArrayRef<ReassociationExprs>":$reassociation,
diff --git a/mlir/include/mlir/Dialect/Tensor/IR/TensorOps.td b/mlir/include/mlir/Dialect/Tensor/IR/TensorOps.td
index a711402609775..beac1e91ee6e3 100644
--- a/mlir/include/mlir/Dialect/Tensor/IR/TensorOps.td
+++ b/mlir/include/mlir/Dialect/Tensor/IR/TensorOps.td
@@ -1218,9 +1218,11 @@ def Tensor_CollapseShapeOp : Tensor_ReassociativeReshapeOp<"collapse_shape"> {
"ArrayRef<ReassociationIndices>":$reassociation,
CArg<"ArrayRef<NamedAttribute>", "{}">:$attrs),
[{
- $_state.addAttribute("reassociation",
- getReassociationIndicesAttribute($_builder, reassociation));
- build($_builder, $_state, resultType, src, attrs);
+ buildPropertiesAndDiscardableAttributes($_state, attrs);
+ $_state.getOrAddProperties<Properties>().reassociation =
+ getReassociationIndicesAttribute($_builder, reassociation);
+ $_state.addOperands(src);
+ $_state.addTypes(resultType);
}]>,
OpBuilder<(ins "Type":$resultType, "Value":$src,
"ArrayRef<ReassociationExprs>":$reassociation,
diff --git a/mlir/include/mlir/TableGen/Operator.h b/mlir/include/mlir/TableGen/Operator.h
index 4c0ba2a1db9ec..bb36c077708e9 100644
--- a/mlir/include/mlir/TableGen/Operator.h
+++ b/mlir/include/mlir/TableGen/Operator.h
@@ -38,6 +38,14 @@ class StringInit;
namespace mlir {
namespace tblgen {
+/// The canonical and legacy names of the implicit segment-size properties.
+inline constexpr StringLiteral operandSegmentAttrName = "operandSegmentSizes";
+inline constexpr StringLiteral resultSegmentAttrName = "resultSegmentSizes";
+inline constexpr StringLiteral legacyOperandSegmentAttrName =
+ "operand_segment_sizes";
+inline constexpr StringLiteral legacyResultSegmentAttrName =
+ "result_segment_sizes";
+
/// This class represents an inferred result type. The result type can be
/// inferred from an argument or result type. If it is inferred from another
/// result type, that type must be buildable or inferred from yet another type.
@@ -209,6 +217,13 @@ class Operator {
llvm::iterator_range<property_iterator> getProperties() { return properties; }
int getNumCoreAttributes() const { return properties.size(); }
+ /// Returns whether this operation has any non-empty properties.
+ bool hasNonEmptyProperties() const;
+
+ /// Returns all accepted attribute spellings for this operation's inherent
+ /// attributes and properties, including legacy segment-size aliases.
+ SmallVector<StringRef> getInherentAttrNames() const;
+
// Op properties accessors.
NamedProperty &getProperty(int index) { return properties[index]; }
const NamedProperty &getProperty(int index) const {
diff --git a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
index 2a846339b562e..efc756d04cbed 100644
--- a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
+++ b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
@@ -497,7 +497,7 @@ struct RawBufferOpLowering : public ConvertOpToLLVMPattern<GpuOp> {
llvm::SmallVector<Type, 1> resultTypes(gpuOp->getNumResults(),
llvmBufferValType);
- typename Intrinsic::Properties properties;
+ typename Intrinsic::Properties properties{};
properties.aux = rewriter.getI32IntegerAttr(0);
Operation *lowered =
Intrinsic::create(rewriter, loc, resultTypes, args, properties);
@@ -2252,33 +2252,38 @@ struct TransposeLoadOpLowering
case 4: {
if (numElements != 16)
return emitNumElementsError(16, "gfx1250+");
- intrinsic =
- ROCDL::DsLoadTr4_B64::create(rewriter, loc, rocdlResultType, srcPtr)
- .getResult();
+ intrinsic = ROCDL::DsLoadTr4_B64::create(
+ rewriter, loc, TypeRange{rocdlResultType},
+ ValueRange{srcPtr}, ROCDL::DsLoadTr4_B64::Properties{})
+ .getResult();
break;
}
case 6: {
if (numElements != 16)
return emitNumElementsError(16, "gfx1250+");
- intrinsic =
- ROCDL::DsLoadTr6_B96::create(rewriter, loc, rocdlResultType, srcPtr)
- .getResult();
+ intrinsic = ROCDL::DsLoadTr6_B96::create(
+ rewriter, loc, TypeRange{rocdlResultType},
+ ValueRange{srcPtr}, ROCDL::DsLoadTr6_B96::Properties{})
+ .getResult();
break;
}
case 8: {
if (numElements != 8)
return emitNumElementsError(8, "gfx1250+");
- intrinsic =
- ROCDL::DsLoadTr8_B64::create(rewriter, loc, rocdlResultType, srcPtr)
- .getResult();
+ intrinsic = ROCDL::DsLoadTr8_B64::create(
+ rewriter, loc, TypeRange{rocdlResultType},
+ ValueRange{srcPtr}, ROCDL::DsLoadTr8_B64::Properties{})
+ .getResult();
break;
}
case 16: {
if (numElements != 8)
return emitNumElementsError(8, "gfx1250+");
- intrinsic = ROCDL::DsLoadTr16_B128::create(rewriter, loc,
- rocdlResultType, srcPtr)
- .getResult();
+ intrinsic =
+ ROCDL::DsLoadTr16_B128::create(
+ rewriter, loc, TypeRange{rocdlResultType}, ValueRange{srcPtr},
+ ROCDL::DsLoadTr16_B128::Properties{})
+ .getResult();
break;
}
default:
@@ -2289,33 +2294,41 @@ struct TransposeLoadOpLowering
case 4: {
if (numElements != 16)
return emitNumElementsError(16, "gfx950");
- intrinsic = ROCDL::ds_read_tr4_b64::create(rewriter, loc,
- rocdlResultType, srcPtr)
- .getResult();
+ intrinsic =
+ ROCDL::ds_read_tr4_b64::create(
+ rewriter, loc, TypeRange{rocdlResultType}, ValueRange{srcPtr},
+ ROCDL::ds_read_tr4_b64::Properties{})
+ .getResult();
break;
}
case 6: {
if (numElements != 16)
return emitNumElementsError(16, "gfx950");
- intrinsic = ROCDL::ds_read_tr6_b96::create(rewriter, loc,
- rocdlResultType, srcPtr)
- .getResult();
+ intrinsic =
+ ROCDL::ds_read_tr6_b96::create(
+ rewriter, loc, TypeRange{rocdlResultType}, ValueRange{srcPtr},
+ ROCDL::ds_read_tr6_b96::Properties{})
+ .getResult();
break;
}
case 8: {
if (numElements != 8)
return emitNumElementsError(8, "gfx950");
- intrinsic = ROCDL::ds_read_tr8_b64::create(rewriter, loc,
- rocdlResultType, srcPtr)
- .getResult();
+ intrinsic =
+ ROCDL::ds_read_tr8_b64::create(
+ rewriter, loc, TypeRange{rocdlResultType}, ValueRange{srcPtr},
+ ROCDL::ds_read_tr8_b64::Properties{})
+ .getResult();
break;
}
case 16: {
if (numElements != 4)
return emitNumElementsError(4, "gfx950");
- intrinsic = ROCDL::ds_read_tr16_b64::create(rewriter, loc,
- rocdlResultType, srcPtr)
- .getResult();
+ intrinsic =
+ ROCDL::ds_read_tr16_b64::create(
+ rewriter, loc, TypeRange{rocdlResultType}, ValueRange{srcPtr},
+ ROCDL::ds_read_tr16_b64::Properties{})
+ .getResult();
break;
}
default:
@@ -2376,8 +2389,9 @@ struct GlobalTransposeLoadOpLowering
assert(numElements == 16);
if (chipset < kGfx1250)
return op.emitOpError("4-bit global_transpose_load requires gfx1250+");
- auto rocdlOp = ROCDL::GlobalLoadTr4_B64::create(rewriter, loc,
- rocdlResultType, srcPtr);
+ auto rocdlOp = ROCDL::GlobalLoadTr4_B64::create(
+ rewriter, loc, TypeRange{rocdlResultType}, ValueRange{srcPtr},
+ ROCDL::GlobalLoadTr4_B64::Properties{});
rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(op, llvmResultType, rocdlOp);
break;
}
@@ -2385,22 +2399,25 @@ struct GlobalTransposeLoadOpLowering
assert(numElements == 16);
if (chipset < kGfx1250)
return op.emitOpError("6-bit global_transpose_load requires gfx1250+");
- auto rocdlOp = ROCDL::GlobalLoadTr6_B96::create(rewriter, loc,
- rocdlResultType, srcPtr);
+ auto rocdlOp = ROCDL::GlobalLoadTr6_B96::create(
+ rewriter, loc, TypeRange{rocdlResultType}, ValueRange{srcPtr},
+ ROCDL::GlobalLoadTr6_B96::Properties{});
rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(op, llvmResultType, rocdlOp);
break;
}
case 8: {
assert(numElements == 8);
- auto rocdlOp = ROCDL::GlobalLoadTr8_B64::create(rewriter, loc,
- rocdlResultType, srcPtr);
+ auto rocdlOp = ROCDL::GlobalLoadTr8_B64::create(
+ rewriter, loc, TypeRange{rocdlResultType}, ValueRange{srcPtr},
+ ROCDL::GlobalLoadTr8_B64::Properties{});
rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(op, llvmResultType, rocdlOp);
break;
}
case 16: {
assert(numElements == 8);
- rewriter.replaceOpWithNewOp<ROCDL::GlobalLoadTr8_B128>(op, llvmResultType,
- srcPtr);
+ rewriter.replaceOpWithNewOp<ROCDL::GlobalLoadTr8_B128>(
+ op, TypeRange{llvmResultType}, ValueRange{srcPtr},
+ ROCDL::GlobalLoadTr8_B128::Properties{});
break;
}
default:
diff --git a/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp b/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp
index cf44c2a0033ac..83680b6a0042e 100644
--- a/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp
+++ b/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp
@@ -112,7 +112,8 @@ static Value castF32To(Type desType, Value f32, Location loc,
if (elementType.getIntOrFloatBitWidth() < 32)
return arith::TruncFOp::create(rewriter, loc, desType, f32);
if (elementType.getIntOrFloatBitWidth() > 32)
- return arith::ExtFOp::create(rewriter, loc, desType, f32);
+ return arith::ExtFOp::create(rewriter, loc, TypeRange{desType},
+ ValueRange{f32}, arith::ExtFOp::Properties{});
llvm_unreachable("The only 32-bit float type is f32");
}
@@ -151,8 +152,9 @@ ExtFOnFloat8RewritePattern::matchAndRewrite(arith::ExtFOp op,
rewriter.createOrFold<vector::BroadcastOp>(loc, outType, zero);
Value scalarIn =
vector::ExtractOp::create(rewriter, loc, in, ArrayRef<int64_t>{});
- Value scalarExt =
- arith::ExtFOp::create(rewriter, loc, outElemType, scalarIn);
+ Value scalarExt = arith::ExtFOp::create(
+ rewriter, loc, TypeRange{outElemType}, ValueRange{scalarIn},
+ arith::ExtFOp::Properties{});
Value result = vector::InsertOp::create(rewriter, loc, scalarExt,
zerodSplat, ArrayRef<int64_t>{});
rewriter.replaceOp(op, result);
@@ -203,7 +205,9 @@ static Value castToF32(Value value, Location loc, PatternRewriter &rewriter) {
if (type.isF32())
return value;
if (type.getIntOrFloatBitWidth() < 32)
- return arith::ExtFOp::create(rewriter, loc, rewriter.getF32Type(), value);
+ return arith::ExtFOp::create(
+ rewriter, loc, TypeRange{rewriter.getF32Type()}, ValueRange{value},
+ arith::ExtFOp::Properties{});
if (type.getIntOrFloatBitWidth() > 32)
return arith::TruncFOp::create(rewriter, loc, rewriter.getF32Type(), value);
llvm_unreachable("The only 32-bit float type is f32");
@@ -472,7 +476,9 @@ ScalingExtFRewritePattern::matchAndRewrite(arith::ScalingExtFOp op,
Type scaleF32Type =
scaleVecType ? VectorType::get(scaleVecType.getShape(), f32) : f32;
if (scaleType.getIntOrFloatBitWidth() < 32)
- scale = arith::ExtFOp::create(rewriter, loc, scaleF32Type, scale);
+ scale =
+ arith::ExtFOp::create(rewriter, loc, TypeRange{scaleF32Type},
+ ValueRange{scale}, arith::ExtFOp::Properties{});
else if (scaleType.getIntOrFloatBitWidth() > 32)
scale = arith::TruncFOp::create(rewriter, loc, scaleF32Type, scale);
@@ -590,7 +596,9 @@ ScalingTruncFRewritePattern::matchAndRewrite(arith::ScalingTruncFOp op,
Type scaleF32Type =
scaleVecType ? VectorType::get(scaleVecType.getShape(), f32) : f32;
if (scaleType.getIntOrFloatBitWidth() < 32)
- scale = arith::ExtFOp::create(rewriter, loc, scaleF32Type, scale);
+ scale =
+ arith::ExtFOp::create(rewriter, loc, TypeRange{scaleF32Type},
+ ValueRange{scale}, arith::ExtFOp::Properties{});
else if (scaleType.getIntOrFloatBitWidth() > 32)
scale = arith::TruncFOp::create(rewriter, loc, scaleF32Type, scale);
diff --git a/mlir/lib/Conversion/ArithToEmitC/ArithToEmitC.cpp b/mlir/lib/Conversion/ArithToEmitC/ArithToEmitC.cpp
index 91dde9018a288..26fa36b937239 100644
--- a/mlir/lib/Conversion/ArithToEmitC/ArithToEmitC.cpp
+++ b/mlir/lib/Conversion/ArithToEmitC/ArithToEmitC.cpp
@@ -708,8 +708,9 @@ class FtoICastOpConversion : public OpConversionPattern<CastOp> {
/*isSigned=*/false);
}
- Value result = emitc::CastOp::create(
- rewriter, castOp.getLoc(), actualResultType, adaptor.getOperands());
+ Value result =
+ emitc::CastOp::create(rewriter, castOp.getLoc(), actualResultType,
+ adaptor.getOperands().front(), /*pure=*/false);
if (isa<arith::FPToUIOp>(castOp)) {
result =
diff --git a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp
index 185c20ab9bb1c..f5e21d63d20b6 100644
--- a/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp
+++ b/mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp
@@ -600,17 +600,19 @@ struct CallOpInterfaceLowering : public ConvertOpToLLVMPattern<CallOpType> {
auto promoted = this->getTypeConverter()->promoteOperands(
callOp.getLoc(), /*opOperands=*/callOp->getOperands(),
adaptor.getOperands(), rewriter, useBarePtrCallConv);
+ LLVM::CallOp::Properties properties{};
+ LLVM::CallOp::populateDefaultProperties(
+ OperationName(LLVM::CallOp::getOperationName(), rewriter.getContext()),
+ properties);
+ properties.operandSegmentSizes = {static_cast<int32_t>(promoted.size()), 0};
+ properties.op_bundle_sizes = rewriter.getDenseI32ArrayAttr({});
auto newOp = LLVM::CallOp::create(
rewriter, callOp.getLoc(),
packedResult ? TypeRange(packedResult) : TypeRange(), promoted,
- callOp->getDiscardableAttrDictionary().getValue());
+ properties, callOp->getDiscardableAttrDictionary().getValue());
if constexpr (std::is_same_v<CallOpType, func::CallOp>)
newOp.setCalleeAttr(callOp.getCalleeAttr());
- newOp.getProperties().operandSegmentSizes = {
- static_cast<int32_t>(promoted.size()), 0};
- newOp.getProperties().op_bundle_sizes = rewriter.getDenseI32ArrayAttr({});
-
// Helper function that extracts an individual result from the return value
// of the new call op. llvm.call ops support only 0 or 1 result. In case of
// 2 or more results, the results are packed into a structure.
diff --git a/mlir/lib/Conversion/MathToFuncs/MathToFuncs.cpp b/mlir/lib/Conversion/MathToFuncs/MathToFuncs.cpp
index 9ff2e3e4fdf7f..f653c6b38e002 100644
--- a/mlir/lib/Conversion/MathToFuncs/MathToFuncs.cpp
+++ b/mlir/lib/Conversion/MathToFuncs/MathToFuncs.cpp
@@ -128,8 +128,9 @@ VecOpToScalarOp<Op>::matchAndRewrite(Op op, PatternRewriter &rewriter) const {
for (Value input : op->getOperands())
operands.push_back(
vector::ExtractOp::create(rewriter, loc, input, positions));
- Value scalarOp =
- Op::create(rewriter, loc, vecType.getElementType(), operands);
+ Value scalarOp = Op::create(
+ rewriter, loc, TypeRange{vecType.getElementType()}, operands,
+ op.getProperties(), op->getDiscardableAttrDictionary().getValue());
result =
vector::InsertOp::create(rewriter, loc, scalarOp, result, positions);
}
diff --git a/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp b/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp
index 1ed1ffd005d49..884a274dfdc73 100644
--- a/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp
+++ b/mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp
@@ -188,7 +188,8 @@ struct SincosOpLowering
rewriter.getContext(), {llvmOperandType, llvmOperandType});
auto sincosOp = LLVM::SincosOp::create(
- rewriter, loc, structType, adaptor.getOperand(), attrs.getAttrs());
+ rewriter, loc, TypeRange{structType}, ValueRange{adaptor.getOperand()},
+ attrs.getProperties(), attrs.getAttrs());
auto sinValue = LLVM::ExtractValueOp::create(rewriter, loc, sincosOp, 0);
auto cosValue = LLVM::ExtractValueOp::create(rewriter, loc, sincosOp, 1);
@@ -233,10 +234,13 @@ struct ExpM1OpLowering
one =
LLVM::ConstantOp::create(rewriter, loc, llvmOperandType, floatOne);
}
- auto exp = LLVM::ExpOp::create(rewriter, loc, adaptor.getOperand(),
- expAttrs.getAttrs());
+ auto exp =
+ LLVM::ExpOp::create(rewriter, loc, TypeRange{llvmOperandType},
+ ValueRange{adaptor.getOperand()},
+ expAttrs.getProperties(), expAttrs.getAttrs());
rewriter.replaceOpWithNewOp<LLVM::FSubOp>(
- op, llvmOperandType, ValueRange{exp, one}, subAttrs.getAttrs());
+ op, TypeRange{llvmOperandType}, ValueRange{exp, one},
+ subAttrs.getProperties(), subAttrs.getAttrs());
return success();
}
@@ -253,11 +257,12 @@ struct ExpM1OpLowering
floatOne);
auto one = LLVM::ConstantOp::create(rewriter, loc, llvm1DVectorTy,
splatAttr);
- auto exp = LLVM::ExpOp::create(rewriter, loc, llvm1DVectorTy,
- operands[0], expAttrs.getAttrs());
- return LLVM::FSubOp::create(rewriter, loc, llvm1DVectorTy,
- ValueRange{exp, one},
- subAttrs.getAttrs());
+ auto exp = LLVM::ExpOp::create(
+ rewriter, loc, TypeRange{llvm1DVectorTy}, ValueRange{operands[0]},
+ expAttrs.getProperties(), expAttrs.getAttrs());
+ return LLVM::FSubOp::create(
+ rewriter, loc, TypeRange{llvm1DVectorTy}, ValueRange{exp, one},
+ subAttrs.getProperties(), subAttrs.getAttrs());
},
rewriter);
}
@@ -297,11 +302,13 @@ struct Log1pOpLowering
: LLVM::ConstantOp::create(rewriter, loc, llvmOperandType,
floatOne);
- auto add = LLVM::FAddOp::create(rewriter, loc, llvmOperandType,
- ValueRange{one, adaptor.getOperand()},
- addAttrs.getAttrs());
+ auto add =
+ LLVM::FAddOp::create(rewriter, loc, TypeRange{llvmOperandType},
+ ValueRange{one, adaptor.getOperand()},
+ addAttrs.getProperties(), addAttrs.getAttrs());
rewriter.replaceOpWithNewOp<LLVM::LogOp>(
- op, llvmOperandType, ValueRange{add}, logAttrs.getAttrs());
+ op, TypeRange{llvmOperandType}, ValueRange{add},
+ logAttrs.getProperties(), logAttrs.getAttrs());
return success();
}
@@ -318,11 +325,13 @@ struct Log1pOpLowering
floatOne);
auto one = LLVM::ConstantOp::create(rewriter, loc, llvm1DVectorTy,
splatAttr);
- auto add = LLVM::FAddOp::create(rewriter, loc, llvm1DVectorTy,
- ValueRange{one, operands[0]},
- addAttrs.getAttrs());
- return LLVM::LogOp::create(rewriter, loc, llvm1DVectorTy,
- ValueRange{add}, logAttrs.getAttrs());
+ auto add = LLVM::FAddOp::create(
+ rewriter, loc, TypeRange{llvm1DVectorTy},
+ ValueRange{one, operands[0]}, addAttrs.getProperties(),
+ addAttrs.getAttrs());
+ return LLVM::LogOp::create(rewriter, loc, TypeRange{llvm1DVectorTy},
+ ValueRange{add}, logAttrs.getProperties(),
+ logAttrs.getAttrs());
},
rewriter);
}
@@ -363,10 +372,13 @@ struct RsqrtOpLowering
one =
LLVM::ConstantOp::create(rewriter, loc, llvmOperandType, floatOne);
}
- auto sqrt = LLVM::SqrtOp::create(rewriter, loc, adaptor.getOperand(),
- sqrtAttrs.getAttrs());
+ auto sqrt =
+ LLVM::SqrtOp::create(rewriter, loc, TypeRange{llvmOperandType},
+ ValueRange{adaptor.getOperand()},
+ sqrtAttrs.getProperties(), sqrtAttrs.getAttrs());
rewriter.replaceOpWithNewOp<LLVM::FDivOp>(
- op, llvmOperandType, ValueRange{one, sqrt}, divAttrs.getAttrs());
+ op, TypeRange{llvmOperandType}, ValueRange{one, sqrt},
+ divAttrs.getProperties(), divAttrs.getAttrs());
return success();
}
@@ -383,11 +395,12 @@ struct RsqrtOpLowering
floatOne);
auto one = LLVM::ConstantOp::create(rewriter, loc, llvm1DVectorTy,
splatAttr);
- auto sqrt = LLVM::SqrtOp::create(rewriter, loc, llvm1DVectorTy,
- operands[0], sqrtAttrs.getAttrs());
- return LLVM::FDivOp::create(rewriter, loc, llvm1DVectorTy,
- ValueRange{one, sqrt},
- divAttrs.getAttrs());
+ auto sqrt = LLVM::SqrtOp::create(
+ rewriter, loc, TypeRange{llvm1DVectorTy}, ValueRange{operands[0]},
+ sqrtAttrs.getProperties(), sqrtAttrs.getAttrs());
+ return LLVM::FDivOp::create(
+ rewriter, loc, TypeRange{llvm1DVectorTy}, ValueRange{one, sqrt},
+ divAttrs.getProperties(), divAttrs.getAttrs());
},
rewriter);
}
diff --git a/mlir/lib/Conversion/MathToLibm/MathToLibm.cpp b/mlir/lib/Conversion/MathToLibm/MathToLibm.cpp
index 1f8020152807a..8039947752b32 100644
--- a/mlir/lib/Conversion/MathToLibm/MathToLibm.cpp
+++ b/mlir/lib/Conversion/MathToLibm/MathToLibm.cpp
@@ -96,8 +96,9 @@ VecOpToScalarOp<Op>::matchAndRewrite(Op op, PatternRewriter &rewriter) const {
for (auto input : op->getOperands())
operands.push_back(
vector::ExtractOp::create(rewriter, loc, input, positions));
- Value scalarOp =
- Op::create(rewriter, loc, vecType.getElementType(), operands);
+ Value scalarOp = Op::create(
+ rewriter, loc, TypeRange{vecType.getElementType()}, operands,
+ op.getProperties(), op->getDiscardableAttrDictionary().getValue());
result =
vector::InsertOp::create(rewriter, loc, scalarOp, result, positions);
}
@@ -116,9 +117,13 @@ PromoteOpToF32<Op>::matchAndRewrite(Op op, PatternRewriter &rewriter) const {
auto f32 = rewriter.getF32Type();
auto extendedOperands =
llvm::map_to_vector(op->getOperands(), [&](Value operand) -> Value {
- return arith::ExtFOp::create(rewriter, loc, f32, operand);
+ return arith::ExtFOp::create(rewriter, loc, TypeRange{f32},
+ ValueRange{operand},
+ arith::ExtFOp::Properties{});
});
- auto newOp = Op::create(rewriter, loc, f32, extendedOperands);
+ auto newOp = Op::create(rewriter, loc, TypeRange{f32}, extendedOperands,
+ op.getProperties(),
+ op->getDiscardableAttrDictionary().getValue());
rewriter.replaceOpWithNewOp<arith::TruncFOp>(op, opType, newOp);
return success();
}
diff --git a/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp b/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
index 7143dcc19ff10..048fb49f6ccdc 100644
--- a/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
+++ b/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
@@ -144,8 +144,8 @@ struct ConvertNativeFuncPattern final : public OpConversionPattern<Op> {
// calls, in order to allow further fastmath optimizations: We thus need to
// convert arith fastmath attrs into attrs recognized by llvm.
arith::AttrConvertFastMathToLLVM<Op, LLVM::CallOp> fastAttrConverter(op);
- mlir::NamedAttribute fastAttr = fastAttrConverter.getAttrs()[0];
- callOp->setAttr(fastAttr.getName(), fastAttr.getValue());
+ callOp.setFastmathFlagsAttr(
+ fastAttrConverter.getProperties().getFastmathFlags());
if (unwrapSizeOneVec) {
// Re-wrap the scalar result back into a size-1 vector to preserve types.
diff --git a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
index 7207c960985f3..3741632717891 100644
--- a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
+++ b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
@@ -863,7 +863,7 @@ struct NVGPUMBarrierArriveLowering
Value barrier =
getMbarrierPtr(b, op.getBarriers().getType(), adaptor.getBarriers(),
adaptor.getMbarId(), rewriter);
- rewriter.replaceOpWithNewOp<NVVM::MBarrierArriveOp>(op, barrier);
+ rewriter.replaceOpWithNewOp<NVVM::MBarrierArriveOp>(op, barrier, Value{});
return success();
}
};
diff --git a/mlir/lib/Conversion/RaiseWasm/RaiseWasmMLIR.cpp b/mlir/lib/Conversion/RaiseWasm/RaiseWasmMLIR.cpp
index c26a3b038e918..fc0a7f1f3851a 100644
--- a/mlir/lib/Conversion/RaiseWasm/RaiseWasmMLIR.cpp
+++ b/mlir/lib/Conversion/RaiseWasm/RaiseWasmMLIR.cpp
@@ -40,6 +40,15 @@ using namespace mlir;
using namespace mlir::wasmssa;
namespace {
+template <typename OpTy>
+static typename OpTy::Properties getDefaultProperties(MLIRContext *context) {
+ typename OpTy::Properties properties{};
+ if constexpr (!std::is_same_v<typename OpTy::Properties, EmptyProperties>)
+ OpTy::populateDefaultProperties(
+ OperationName(OpTy::getOperationName(), context), properties);
+ return properties;
+}
+
template <typename SourceOp, typename TargetIntOp, typename TargetFPOp>
struct IntFPDispatchMappingConversion : OpConversionPattern<SourceOp> {
using OpConversionPattern<SourceOp>::OpConversionPattern;
@@ -49,14 +58,18 @@ struct IntFPDispatchMappingConversion : OpConversionPattern<SourceOp> {
ConversionPatternRewriter &rewriter) const override {
Type type = srcOp.getRhs().getType();
if (type.isInteger()) {
- rewriter.replaceOpWithNewOp<TargetIntOp>(srcOp, srcOp->getResultTypes(),
- adaptor.getOperands());
+ rewriter.replaceOpWithNewOp<TargetIntOp>(
+ srcOp, srcOp->getResultTypes(), adaptor.getOperands(),
+ getDefaultProperties<TargetIntOp>(rewriter.getContext()),
+ ArrayRef<NamedAttribute>{});
return success();
}
if (!type.isFloat())
return failure();
- rewriter.replaceOpWithNewOp<TargetFPOp>(srcOp, srcOp->getResultTypes(),
- adaptor.getOperands());
+ rewriter.replaceOpWithNewOp<TargetFPOp>(
+ srcOp, srcOp->getResultTypes(), adaptor.getOperands(),
+ getDefaultProperties<TargetFPOp>(rewriter.getContext()),
+ ArrayRef<NamedAttribute>{});
return success();
}
};
@@ -77,8 +90,10 @@ struct OpMappingConversion : OpConversionPattern<SourceOp> {
LogicalResult
matchAndRewrite(SourceOp srcOp, typename SourceOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- rewriter.replaceOpWithNewOp<TargetOp>(srcOp, srcOp->getResultTypes(),
- adaptor.getOperands());
+ rewriter.replaceOpWithNewOp<TargetOp>(
+ srcOp, srcOp->getResultTypes(), adaptor.getOperands(),
+ getDefaultProperties<TargetOp>(rewriter.getContext()),
+ ArrayRef<NamedAttribute>{});
return success();
}
};
diff --git a/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp b/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp
index 370457c85e797..abbae7466ad9f 100644
--- a/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp
+++ b/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp
@@ -714,7 +714,8 @@ ParallelToGpuLaunchLowering::matchAndRewrite(ParallelOp parallelOp,
if (!externalValues.empty())
return failure();
// Replace by gpu.all_reduce.
- auto gpuRedOp = gpu::AllReduceOp::create(rewriter, loc, newValue);
+ auto gpuRedOp = gpu::AllReduceOp::create(
+ rewriter, loc, newValue, /*op=*/nullptr, /*uniform=*/false);
cloningMap.map(parentLoop->getResult(0), gpuRedOp.getResult());
// Copy region.
rewriter.inlineRegionBefore(reduceOp.getRegion(0), gpuRedOp.getRegion(),
diff --git a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
index 04c2288b71c2c..4e733ff45360d 100644
--- a/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
+++ b/mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
@@ -303,13 +303,45 @@ static Type convertStructType(spirv::StructType type,
namespace {
-template <typename OpTy>
-static NamedAttrList collectAttrsForConversion(OpTy op) {
+template <typename TargetOp, typename SourceOp>
+static LogicalResult
+collectAttrsForConversion(SourceOp op,
+ typename TargetOp::Properties &properties,
+ SmallVectorImpl<NamedAttribute> &discardableAttrs) {
NamedAttrList attrs(op->getDiscardableAttrDictionary());
- if (auto properties =
+ if (auto sourceProperties =
dyn_cast_or_null<DictionaryAttr>(op->getPropertiesAsAttribute()))
- attrs.append(properties.getValue());
- return attrs;
+ attrs.append(sourceProperties.getValue());
+
+ if constexpr (std::is_same_v<typename TargetOp::Properties,
+ EmptyProperties>) {
+ llvm::append_range(discardableAttrs, attrs);
+ return success();
+ }
+
+ TargetOp::populateDefaultProperties(
+ OperationName(TargetOp::getOperationName(), op->getContext()),
+ properties);
+ if (failed(TargetOp::setPropertiesFromAttr(
+ properties, attrs.getDictionary(op->getContext()),
+ [&]() { return op.emitError("failed to convert properties"); })))
+ return failure();
+
+ auto convertedProperties = dyn_cast_or_null<DictionaryAttr>(
+ TargetOp::getPropertiesAsAttr(op->getContext(), properties));
+ for (NamedAttribute attr : attrs) {
+ StringRef name = attr.getName().getValue();
+ bool isProperty = convertedProperties && convertedProperties.contains(name);
+ if (name == "operand_segment_sizes")
+ isProperty |= convertedProperties &&
+ convertedProperties.contains("operandSegmentSizes");
+ if (name == "result_segment_sizes")
+ isProperty |= convertedProperties &&
+ convertedProperties.contains("resultSegmentSizes");
+ if (!isProperty)
+ discardableAttrs.push_back(attr);
+ }
+ return success();
}
class AccessChainPattern : public SPIRVToLLVMConversion<spirv::AccessChainOp> {
@@ -442,9 +474,13 @@ class ConstantScalarAndVectorPattern
rewriter.replaceOpWithNewOp<LLVM::ConstantOp>(constOp, dstType, dstAttr);
return success();
}
+ LLVM::ConstantOp::Properties properties{};
+ SmallVector<NamedAttribute> discardableAttrs;
+ if (failed(collectAttrsForConversion<LLVM::ConstantOp>(constOp, properties,
+ discardableAttrs)))
+ return failure();
rewriter.replaceOpWithNewOp<LLVM::ConstantOp>(
- constOp, dstType, adaptor.getOperands(),
- collectAttrsForConversion(constOp));
+ constOp, dstType, adaptor.getOperands(), properties, discardableAttrs);
return success();
}
};
@@ -653,8 +689,13 @@ class DirectConversionPattern : public SPIRVToLLVMConversion<SPIRVOp> {
auto dstType = this->getTypeConverter()->convertType(op.getType());
if (!dstType)
return rewriter.notifyMatchFailure(op, "type conversion failed");
+ typename LLVMOp::Properties properties{};
+ SmallVector<NamedAttribute> discardableAttrs;
+ if (failed(collectAttrsForConversion<LLVMOp>(op, properties,
+ discardableAttrs)))
+ return failure();
rewriter.template replaceOpWithNewOp<LLVMOp>(
- op, dstType, adaptor.getOperands(), collectAttrsForConversion(op));
+ op, dstType, adaptor.getOperands(), properties, discardableAttrs);
return success();
}
};
@@ -887,13 +928,19 @@ class FunctionCallPattern
LogicalResult
matchAndRewrite(spirv::FunctionCallOp callOp, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
+ LLVM::CallOp::Properties properties{};
+ SmallVector<NamedAttribute> discardableAttrs;
+ if (failed(collectAttrsForConversion<LLVM::CallOp>(callOp, properties,
+ discardableAttrs)))
+ return failure();
+ properties.operandSegmentSizes = {
+ static_cast<int32_t>(adaptor.getOperands().size()), 0};
+ properties.op_bundle_sizes = rewriter.getDenseI32ArrayAttr({});
+
if (callOp.getNumResults() == 0) {
- auto newOp = rewriter.replaceOpWithNewOp<LLVM::CallOp>(
- callOp, TypeRange(), adaptor.getOperands(),
- collectAttrsForConversion(callOp));
- newOp.getProperties().operandSegmentSizes = {
- static_cast<int32_t>(adaptor.getOperands().size()), 0};
- newOp.getProperties().op_bundle_sizes = rewriter.getDenseI32ArrayAttr({});
+ rewriter.replaceOpWithNewOp<LLVM::CallOp>(callOp, TypeRange(),
+ adaptor.getOperands(),
+ properties, discardableAttrs);
return success();
}
@@ -901,12 +948,8 @@ class FunctionCallPattern
auto dstType = getTypeConverter()->convertType(callOp.getType(0));
if (!dstType)
return rewriter.notifyMatchFailure(callOp, "type conversion failed");
- auto newOp = rewriter.replaceOpWithNewOp<LLVM::CallOp>(
- callOp, dstType, adaptor.getOperands(),
- collectAttrsForConversion(callOp));
- newOp.getProperties().operandSegmentSizes = {
- static_cast<int32_t>(adaptor.getOperands().size()), 0};
- newOp.getProperties().op_bundle_sizes = rewriter.getDenseI32ArrayAttr({});
+ rewriter.replaceOpWithNewOp<LLVM::CallOp>(
+ callOp, dstType, adaptor.getOperands(), properties, discardableAttrs);
return success();
}
};
@@ -1951,9 +1994,14 @@ class BitcastConversionPattern
return success();
}
- rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(
- bitcastOp, dstType, adaptor.getOperands(),
- collectAttrsForConversion(bitcastOp));
+ LLVM::BitcastOp::Properties properties{};
+ SmallVector<NamedAttribute> discardableAttrs;
+ if (failed(collectAttrsForConversion<LLVM::BitcastOp>(bitcastOp, properties,
+ discardableAttrs)))
+ return failure();
+ rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(bitcastOp, dstType,
+ adaptor.getOperands(),
+ properties, discardableAttrs);
return success();
}
};
diff --git a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp
index b7eb0a3aed546..cc8c074489cb5 100644
--- a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp
+++ b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp
@@ -35,6 +35,18 @@
using namespace mlir;
using namespace mlir::tosa;
+template <typename OpTy>
+static OpTy createWithDefaultProperties(OpBuilder &builder, Location loc,
+ TypeRange resultTypes,
+ ValueRange operands) {
+ typename OpTy::Properties properties{};
+ OpTy::populateDefaultProperties(
+ OperationName(OpTy::getOperationName(), builder.getContext()),
+ properties);
+ return OpTy::create(builder, loc, resultTypes, operands, properties,
+ /*discardableAttributes=*/{});
+}
+
// Helper function to materialize the semantically correct compare and select
// operations given a binary operation with a specific NaN propagation mode.
//
@@ -89,7 +101,8 @@ static Value createLinalgBodyCalculationForElementwiseOp(
// tosa::AbsOp
if (isa<tosa::AbsOp>(op) && isa<FloatType>(elementTy))
- return math::AbsFOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<math::AbsFOp>(rewriter, loc, resultTypes,
+ args);
if (isa<tosa::AbsOp>(op) && isa<IntegerType>(elementTy)) {
auto zero = arith::ConstantOp::create(rewriter, loc,
@@ -100,21 +113,26 @@ static Value createLinalgBodyCalculationForElementwiseOp(
// tosa::AddOp
if (isa<tosa::AddOp>(op) && isa<FloatType>(elementTy))
- return arith::AddFOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<arith::AddFOp>(rewriter, loc,
+ resultTypes, args);
if (isa<tosa::AddOp>(op) && isa<IntegerType>(elementTy))
- return arith::AddIOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<arith::AddIOp>(rewriter, loc,
+ resultTypes, args);
// tosa::SubOp
if (isa<tosa::SubOp>(op) && isa<FloatType>(elementTy))
- return arith::SubFOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<arith::SubFOp>(rewriter, loc,
+ resultTypes, args);
if (isa<tosa::SubOp>(op) && isa<IntegerType>(elementTy))
- return arith::SubIOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<arith::SubIOp>(rewriter, loc,
+ resultTypes, args);
// tosa::IntDivOp
if (isa<tosa::IntDivOp>(op) && isa<IntegerType>(elementTy))
- return arith::DivSIOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<arith::DivSIOp>(rewriter, loc,
+ resultTypes, args);
// tosa::ReciprocalOp
if (isa<tosa::ReciprocalOp>(op) && isa<FloatType>(elementTy)) {
@@ -282,15 +300,18 @@ static Value createLinalgBodyCalculationForElementwiseOp(
// tosa::LogicalLeftShiftOp
if (isa<tosa::LogicalLeftShiftOp>(op) && isa<IntegerType>(elementTy))
- return arith::ShLIOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<arith::ShLIOp>(rewriter, loc,
+ resultTypes, args);
// tosa::LogicalRightShiftOp
if (isa<tosa::LogicalRightShiftOp>(op) && isa<IntegerType>(elementTy))
- return arith::ShRUIOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<arith::ShRUIOp>(rewriter, loc,
+ resultTypes, args);
// tosa::ArithmeticRightShiftOp
if (isa<tosa::ArithmeticRightShiftOp>(op) && isa<IntegerType>(elementTy)) {
- auto result = arith::ShRSIOp::create(rewriter, loc, resultTypes, args);
+ auto result = createWithDefaultProperties<arith::ShRSIOp>(
+ rewriter, loc, resultTypes, args);
bool round = cast<tosa::ArithmeticRightShiftOp>(op).getRound();
if (!round) {
return result;
@@ -316,8 +337,8 @@ static Value createLinalgBodyCalculationForElementwiseOp(
auto shifted =
arith::ShRSIOp::create(rewriter, loc, resultTypes, args[0], subtract)
->getResults();
- auto truncated = arith::TruncIOp::create(rewriter, loc, i1Ty, shifted,
- ArrayRef<NamedAttribute>());
+ auto truncated = createWithDefaultProperties<arith::TruncIOp>(
+ rewriter, loc, TypeRange{i1Ty}, shifted);
auto isInputOdd =
arith::AndIOp::create(rewriter, loc, i1Ty, truncated, i1one);
// shifted, truncated, isInputOdd can be poison when input2 is 0.
@@ -354,35 +375,43 @@ static Value createLinalgBodyCalculationForElementwiseOp(
// tosa::PowOp
if (isa<tosa::PowOp>(op) && isa<FloatType>(elementTy))
- return mlir::math::PowFOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<mlir::math::PowFOp>(rewriter, loc,
+ resultTypes, args);
// tosa::RsqrtOp
if (isa<tosa::RsqrtOp>(op) && isa<FloatType>(elementTy))
- return mlir::math::RsqrtOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<mlir::math::RsqrtOp>(rewriter, loc,
+ resultTypes, args);
// tosa::LogOp
if (isa<tosa::LogOp>(op) && isa<FloatType>(elementTy))
- return mlir::math::LogOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<mlir::math::LogOp>(rewriter, loc,
+ resultTypes, args);
// tosa::ExpOp
if (isa<tosa::ExpOp>(op) && isa<FloatType>(elementTy))
- return mlir::math::ExpOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<mlir::math::ExpOp>(rewriter, loc,
+ resultTypes, args);
// tosa::SinOp
if (isa<tosa::SinOp>(op) && isa<FloatType>(elementTy))
- return mlir::math::SinOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<mlir::math::SinOp>(rewriter, loc,
+ resultTypes, args);
// tosa::CosOp
if (isa<tosa::CosOp>(op) && isa<FloatType>(elementTy))
- return mlir::math::CosOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<mlir::math::CosOp>(rewriter, loc,
+ resultTypes, args);
// tosa::TanhOp
if (isa<tosa::TanhOp>(op) && isa<FloatType>(elementTy))
- return mlir::math::TanhOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<mlir::math::TanhOp>(rewriter, loc,
+ resultTypes, args);
// tosa::ErfOp
if (isa<tosa::ErfOp>(op) && llvm::isa<FloatType>(elementTy))
- return mlir::math::ErfOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<mlir::math::ErfOp>(rewriter, loc,
+ resultTypes, args);
// tosa::GreaterOp
if (isa<tosa::GreaterOp>(op) && isa<FloatType>(elementTy))
@@ -442,11 +471,13 @@ static Value createLinalgBodyCalculationForElementwiseOp(
// tosa::CeilOp
if (isa<tosa::CeilOp>(op) && isa<FloatType>(elementTy))
- return math::CeilOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<math::CeilOp>(rewriter, loc, resultTypes,
+ args);
// tosa::FloorOp
if (isa<tosa::FloorOp>(op) && isa<FloatType>(elementTy))
- return math::FloorOp::create(rewriter, loc, resultTypes, args);
+ return createWithDefaultProperties<math::FloorOp>(rewriter, loc,
+ resultTypes, args);
// tosa::ClampOp
if (isa<tosa::ClampOp>(op) && isa<FloatType>(elementTy)) {
@@ -558,21 +589,21 @@ static Value createLinalgBodyCalculationForElementwiseOp(
return args.front();
if (isa<FloatType>(srcTy) && isa<FloatType>(dstTy) && bitExtend)
- return arith::ExtFOp::create(rewriter, loc, resultTypes, args,
- ArrayRef<NamedAttribute>());
+ return createWithDefaultProperties<arith::ExtFOp>(rewriter, loc,
+ resultTypes, args);
if (isa<FloatType>(srcTy) && isa<FloatType>(dstTy) && !bitExtend)
- return arith::TruncFOp::create(rewriter, loc, resultTypes, args,
- ArrayRef<NamedAttribute>());
+ return createWithDefaultProperties<arith::TruncFOp>(rewriter, loc,
+ resultTypes, args);
// 1-bit integers need to be treated as signless.
if (srcTy.isInteger(1) && arith::UIToFPOp::areCastCompatible(srcTy, dstTy))
- return arith::UIToFPOp::create(rewriter, loc, resultTypes, args,
- ArrayRef<NamedAttribute>());
+ return createWithDefaultProperties<arith::UIToFPOp>(rewriter, loc,
+ resultTypes, args);
if (srcTy.isInteger(1) && isa<IntegerType>(dstTy) && bitExtend)
- return arith::ExtUIOp::create(rewriter, loc, resultTypes, args,
- ArrayRef<NamedAttribute>());
+ return createWithDefaultProperties<arith::ExtUIOp>(rewriter, loc,
+ resultTypes, args);
// Unsigned integers need an unrealized cast so that they can be passed
// to UIToFP.
@@ -1116,19 +1147,23 @@ static Value createLinalgBodyCalculationForReduceOp(Operation *op,
PatternRewriter &rewriter) {
Location loc = op->getLoc();
if (isa<tosa::ReduceSumOp>(op) && isa<FloatType>(elementTy)) {
- return arith::AddFOp::create(rewriter, loc, args);
+ return createWithDefaultProperties<arith::AddFOp>(
+ rewriter, loc, TypeRange{elementTy}, args);
}
if (isa<tosa::ReduceSumOp>(op) && isa<IntegerType>(elementTy)) {
- return arith::AddIOp::create(rewriter, loc, args);
+ return createWithDefaultProperties<arith::AddIOp>(
+ rewriter, loc, TypeRange{elementTy}, args);
}
if (isa<tosa::ReduceProductOp>(op) && isa<FloatType>(elementTy)) {
- return arith::MulFOp::create(rewriter, loc, args);
+ return createWithDefaultProperties<arith::MulFOp>(
+ rewriter, loc, TypeRange{elementTy}, args);
}
if (isa<tosa::ReduceProductOp>(op) && isa<IntegerType>(elementTy)) {
- return arith::MulIOp::create(rewriter, loc, args);
+ return createWithDefaultProperties<arith::MulIOp>(
+ rewriter, loc, TypeRange{elementTy}, args);
}
if (isa<tosa::ReduceMinOp>(op) && isa<FloatType>(elementTy)) {
@@ -1251,8 +1286,9 @@ static LogicalResult reduceMatchAndRewriteHelper(OpTy op, uint64_t axis,
// If reduction type differs then extend (applicable to reduce_sum)
if (binaryArgs[0].getType() != accTy)
- binaryArgs[0] = arith::ExtFOp::create(nestedBuilder, nestedLoc, accTy,
- binaryArgs[0]);
+ binaryArgs[0] = arith::ExtFOp::create(
+ nestedBuilder, nestedLoc, TypeRange{accTy},
+ ValueRange{binaryArgs[0]}, arith::ExtFOp::Properties{});
auto result = createLinalgBodyCalculationForReduceOp(op, binaryArgs,
accTy, rewriter);
diff --git a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp
index ee57a5a781c5f..4a11a927090bf 100644
--- a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp
+++ b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp
@@ -135,7 +135,9 @@ static mlir::Value linalgBroadcastAndMaybeExt(PatternRewriter &rewriter,
if (resType != biasVal.getType()) {
biasVal =
resultTy.getElementType().isFloat()
- ? arith::ExtFOp::create(builder, loc, resType, biasVal)
+ ? arith::ExtFOp::create(
+ builder, loc, TypeRange{resType},
+ ValueRange{biasVal}, arith::ExtFOp::Properties{})
.getResult()
: arith::ExtSIOp::create(builder, loc, resType,
biasVal)
diff --git a/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp b/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp
index 975fe28399609..5a55899fcd526 100644
--- a/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp
+++ b/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp
@@ -529,7 +529,9 @@ struct CombineTransferReadOpTranspose final
result = arith::ExtUIOp::create(rewriter, loc, op.getType(), result)
.getResult();
else
- result = arith::ExtFOp::create(rewriter, loc, op.getType(), result)
+ result = arith::ExtFOp::create(rewriter, loc, TypeRange{op.getType()},
+ ValueRange{result},
+ arith::ExtFOp::Properties{})
.getResult();
}
diff --git a/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp b/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp
index c9eba6962e6a4..55c1b122f5429 100644
--- a/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp
+++ b/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp
@@ -615,7 +615,8 @@ struct PrepareTransferReadConversion
Location loc = xferOp.getLoc();
memref::StoreOp::create(rewriter, loc, newXfer->getResult(0),
buffers.dataBuffer);
- rewriter.replaceOpWithNewOp<memref::LoadOp>(xferOp, buffers.dataBuffer);
+ rewriter.replaceOpWithNewOp<memref::LoadOp>(xferOp, buffers.dataBuffer,
+ ValueRange{});
return success();
}
@@ -658,7 +659,8 @@ struct PrepareTransferWriteConversion
auto buffers = allocBuffers(rewriter, xferOp);
memref::StoreOp::create(rewriter, loc, xferOp.getVector(),
buffers.dataBuffer);
- auto loadedVec = memref::LoadOp::create(rewriter, loc, buffers.dataBuffer);
+ auto loadedVec =
+ memref::LoadOp::create(rewriter, loc, buffers.dataBuffer, ValueRange{});
rewriter.modifyOpInPlace(xferOp, [&]() {
xferOp.getValueToStoreMutable().assign(loadedVec);
xferOp->setAttr(kPassLabel, rewriter.getUnitAttr());
diff --git a/mlir/lib/Conversion/VectorToXeGPU/VectorToXeGPU.cpp b/mlir/lib/Conversion/VectorToXeGPU/VectorToXeGPU.cpp
index 9863206f14fe1..54ba2e0bcab11 100644
--- a/mlir/lib/Conversion/VectorToXeGPU/VectorToXeGPU.cpp
+++ b/mlir/lib/Conversion/VectorToXeGPU/VectorToXeGPU.cpp
@@ -965,9 +965,9 @@ struct ContractionLowering : public OpRewritePattern<vector::ContractionOp> {
return rewriter.notifyMatchFailure(contractOp,
"Expects operands of rank 4 or less");
- auto dpasOp = xegpu::DpasOp::create(rewriter, loc,
- TypeRange{contractOp.getResultType()},
- ValueRange{lhs, rhs, acc});
+ auto dpasOp = xegpu::DpasOp::create(
+ rewriter, loc, contractOp.getResultType(), lhs, rhs, acc,
+ /*layout_a=*/nullptr, /*layout_b=*/nullptr, /*layout_cd=*/nullptr);
rewriter.replaceOp(contractOp, dpasOp);
return success();
diff --git a/mlir/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp b/mlir/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp
index d6ff8504582c8..c0cc18e53570d 100644
--- a/mlir/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp
+++ b/mlir/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp
@@ -925,8 +925,8 @@ class LoadStoreMatrixToXeVMPattern : public OpConversionPattern<OpType> {
VectorType::get(valOrResVecTy.getShape(), intElemTy);
if constexpr (std::is_same_v<OpType, xegpu::LoadMatrixOp>) {
- Value loadOp =
- xevm::BlockLoadOp::create(rewriter, loc, intVecTy, basePtrLLVM);
+ Value loadOp = xevm::BlockLoadOp::create(
+ rewriter, loc, intVecTy, basePtrLLVM, /*cache_control=*/nullptr);
if (intVecTy != valOrResVecTy) {
loadOp =
vector::BitCastOp::create(rewriter, loc, valOrResVecTy, loadOp);
diff --git a/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp b/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
index 332ac8cdb60c0..2fcf0a41e8ad6 100644
--- a/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
+++ b/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
@@ -87,6 +87,22 @@ static void patchOperandSegmentSizes(ArrayRef<NamedAttribute> attrs,
}
}
+template <typename OpTy>
+static typename OpTy::Properties
+getPropertiesFromAttrs(OpBuilder &builder, Location loc,
+ ArrayRef<NamedAttribute> attrs) {
+ typename OpTy::Properties properties{};
+ OpTy::populateDefaultProperties(
+ OperationName(OpTy::getOperationName(), builder.getContext()),
+ properties);
+ LogicalResult result =
+ OpTy::setPropertiesFromAttr(properties, builder.getDictionaryAttr(attrs),
+ [&]() { return emitError(loc); });
+ assert(succeeded(result) && "failed to convert operation properties");
+ (void)result;
+ return properties;
+}
+
// A helper function to flatten a vector value to a scalar containing its bits,
// returning the value itself if othetwise.
static Value flattenVecToBits(ConversionPatternRewriter &rewriter, Location loc,
@@ -110,16 +126,22 @@ LogicalResult RawBufferAtomicByCasPattern<AtomicOp, ArithOp>::matchAndRewrite(
ConversionPatternRewriter &rewriter) const {
Location loc = atomicOp.getLoc();
- ArrayRef<NamedAttribute> origAttrs = atomicOp->getAttrs();
+ NamedAttrList origProperties;
+ atomicOp->getName().populateInherentAttrs(atomicOp, origProperties);
+ ArrayRef<NamedAttribute> discardableAttrs =
+ atomicOp->getDiscardableAttrDictionary().getValue();
ValueRange operands = adaptor.getOperands();
Value data = operands.take_front()[0];
ValueRange invariantArgs = operands.drop_front();
Type dataType = data.getType();
SmallVector<NamedAttribute> loadAttrs;
- patchOperandSegmentSizes(origAttrs, loadAttrs, DataArgAction::Drop);
- Value initialLoad = RawBufferLoadOp::create(rewriter, loc, dataType,
- invariantArgs, loadAttrs);
+ patchOperandSegmentSizes(origProperties, loadAttrs, DataArgAction::Drop);
+ auto loadProperties =
+ getPropertiesFromAttrs<RawBufferLoadOp>(rewriter, loc, loadAttrs);
+ Value initialLoad =
+ RawBufferLoadOp::create(rewriter, loc, TypeRange{dataType}, invariantArgs,
+ loadProperties, discardableAttrs);
Block *currentBlock = rewriter.getInsertionBlock();
Block *afterAtomic =
rewriter.splitBlock(currentBlock, rewriter.getInsertionPoint());
@@ -135,11 +157,15 @@ LogicalResult RawBufferAtomicByCasPattern<AtomicOp, ArithOp>::matchAndRewrite(
dataType = operated.getType();
SmallVector<NamedAttribute> cmpswapAttrs;
- patchOperandSegmentSizes(origAttrs, cmpswapAttrs, DataArgAction::Duplicate);
+ patchOperandSegmentSizes(origProperties, cmpswapAttrs,
+ DataArgAction::Duplicate);
SmallVector<Value> cmpswapArgs = {operated, prevLoad};
cmpswapArgs.append(invariantArgs.begin(), invariantArgs.end());
- Value atomicRes = RawBufferAtomicCmpswapOp::create(rewriter, loc, dataType,
- cmpswapArgs, cmpswapAttrs);
+ auto cmpswapProperties = getPropertiesFromAttrs<RawBufferAtomicCmpswapOp>(
+ rewriter, loc, cmpswapAttrs);
+ Value atomicRes = RawBufferAtomicCmpswapOp::create(
+ rewriter, loc, TypeRange{dataType}, cmpswapArgs, cmpswapProperties,
+ discardableAttrs);
// We care about exact bitwise equality here, so do some bitcasts.
// These will fold away during lowering to the ROCDL dialect, where
diff --git a/mlir/lib/Dialect/Arith/IR/ArithOps.cpp b/mlir/lib/Dialect/Arith/IR/ArithOps.cpp
index 7157e29800db0..3575e7fe71e34 100644
--- a/mlir/lib/Dialect/Arith/IR/ArithOps.cpp
+++ b/mlir/lib/Dialect/Arith/IR/ArithOps.cpp
@@ -1461,9 +1461,10 @@ struct NarrowExtremum final : OpRewritePattern<TruncOp> {
return failure();
}
- rewriter.replaceOpWithNewOp<ExtremumOp>(truncOp, TypeRange{narrowType},
- ValueRange{lhs, rhs},
- extremumOp->getAttrs());
+ rewriter.replaceOpWithNewOp<ExtremumOp>(
+ truncOp, TypeRange{narrowType}, ValueRange{lhs, rhs},
+ extremumOp.getProperties(),
+ extremumOp->getDiscardableAttrDictionary().getValue());
return success();
}
};
diff --git a/mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp b/mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp
index b6e101952676a..a509f8b8755e0 100644
--- a/mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp
+++ b/mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp
@@ -98,7 +98,8 @@ void mlir::arith::populateEmulateUnsupportedFloatsConversions(
});
converter.addTargetMaterialization(
[](OpBuilder &b, Type target, ValueRange input, Location loc) {
- auto extFOp = arith::ExtFOp::create(b, loc, target, input);
+ auto extFOp = arith::ExtFOp::create(b, loc, target, input.front(),
+ arith::FastMathFlagsAttr{});
extFOp.setFastmath(arith::FastMathFlags::contract);
return extFOp;
});
diff --git a/mlir/lib/Dialect/Arith/Transforms/ExpandOps.cpp b/mlir/lib/Dialect/Arith/Transforms/ExpandOps.cpp
index 02ed6ccd87a42..1d481c393b374 100644
--- a/mlir/lib/Dialect/Arith/Transforms/ExpandOps.cpp
+++ b/mlir/lib/Dialect/Arith/Transforms/ExpandOps.cpp
@@ -542,7 +542,8 @@ struct F4E2M1TruncFOpConverter : public OpRewritePattern<arith::TruncFOp> {
if (!isa<Float4E2M1FNType>(resultETy))
return rewriter.notifyMatchFailure(op, "not a trunc of F4E2M1FN");
if (!isa<Float32Type>(operandETy))
- operand = arith::ExtFOp::create(b, f32Ty, operand);
+ operand =
+ arith::ExtFOp::create(b, f32Ty, operand, arith::FastMathFlagsAttr{});
Value c0x1 = createConst(loc, i4Ty, 1, rewriter);
Value c0x3 = createConst(loc, i4Ty, 3, rewriter);
diff --git a/mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp b/mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp
index c9eaa66d6ea49..a0b8b53c1ef34 100644
--- a/mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp
+++ b/mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp
@@ -92,8 +92,13 @@ struct ConvertOpToUnsigned final : OpRewritePattern<Signed> {
staticallyNonNegative(this->solver, static_cast<Operation *>(op))))
return failure();
- rw.replaceOpWithNewOp<Unsigned>(op, op->getResultTypes(), op->getOperands(),
- op->getAttrs());
+ typename Unsigned::Properties properties{};
+ Unsigned::populateDefaultProperties(
+ OperationName(Unsigned::getOperationName(), rw.getContext()),
+ properties);
+ rw.replaceOpWithNewOp<Unsigned>(
+ op, op->getResultTypes(), op->getOperands(), properties,
+ op->getDiscardableAttrDictionary().getValue());
return success();
}
diff --git a/mlir/lib/Dialect/Arith/Utils/Utils.cpp b/mlir/lib/Dialect/Arith/Utils/Utils.cpp
index 106d125de14b0..990fb2cd2bde0 100644
--- a/mlir/lib/Dialect/Arith/Utils/Utils.cpp
+++ b/mlir/lib/Dialect/Arith/Utils/Utils.cpp
@@ -177,7 +177,8 @@ static Value convertScalarToFpDtype(ImplicitLocOpBuilder &b, Value operand,
}
if (auto fromFpTy = dyn_cast<FloatType>(operand.getType())) {
if (toType.getWidth() > fromFpTy.getWidth())
- return arith::ExtFOp::create(b, toType, operand);
+ return arith::ExtFOp::create(b, toType, operand,
+ arith::FastMathFlagsAttr{});
if (toType.getWidth() < fromFpTy.getWidth())
return arith::TruncFOp::create(b, toType, operand);
return operand;
@@ -200,8 +201,10 @@ static Value convertScalarToComplexDtype(ImplicitLocOpBuilder &b, Value operand,
real = arith::TruncFOp::create(b, targetETy, real);
imag = arith::TruncFOp::create(b, targetETy, imag);
} else {
- real = arith::ExtFOp::create(b, targetETy, real);
- imag = arith::ExtFOp::create(b, targetETy, imag);
+ real = arith::ExtFOp::create(b, targetETy, real,
+ arith::FastMathFlagsAttr{});
+ imag = arith::ExtFOp::create(b, targetETy, imag,
+ arith::FastMathFlagsAttr{});
}
return complex::CreateOp::create(b, targetType, real, imag);
}
@@ -212,7 +215,7 @@ static Value convertScalarToComplexDtype(ImplicitLocOpBuilder &b, Value operand,
auto toBitwidth = toFpTy.getIntOrFloatBitWidth();
Value from = operand;
if (from.getType().getIntOrFloatBitWidth() < toBitwidth) {
- from = arith::ExtFOp::create(b, toFpTy, from);
+ from = arith::ExtFOp::create(b, toFpTy, from, arith::FastMathFlagsAttr{});
}
if (from.getType().getIntOrFloatBitWidth() > toBitwidth) {
from = arith::TruncFOp::create(b, toFpTy, from);
diff --git a/mlir/lib/Dialect/GPU/Transforms/DecomposeMemRefs.cpp b/mlir/lib/Dialect/GPU/Transforms/DecomposeMemRefs.cpp
index 7b30906abc2fd..ce9237eb1aa14 100644
--- a/mlir/lib/Dialect/GPU/Transforms/DecomposeMemRefs.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/DecomposeMemRefs.cpp
@@ -142,7 +142,7 @@ struct FlattenLoad : public OpRewritePattern<memref::LoadOp> {
Location loc = op.getLoc();
Value flatMemref = getFlatMemref(rewriter, loc, memref, op.getIndices());
- rewriter.replaceOpWithNewOp<memref::LoadOp>(op, flatMemref);
+ rewriter.replaceOpWithNewOp<memref::LoadOp>(op, flatMemref, ValueRange{});
return success();
}
};
diff --git a/mlir/lib/Dialect/Linalg/Transforms/Vectorization.cpp b/mlir/lib/Dialect/Linalg/Transforms/Vectorization.cpp
index e8f8d737e86de..087ef7be2d077 100644
--- a/mlir/lib/Dialect/Linalg/Transforms/Vectorization.cpp
+++ b/mlir/lib/Dialect/Linalg/Transforms/Vectorization.cpp
@@ -3679,7 +3679,8 @@ struct Conv1DGenerator
if (isa<FloatType>(srcElementType) && isa<FloatType>(dstElementType) &&
srcWidth < dstWidth)
- return arith::ExtFOp::create(rewriter, loc, dstType, val);
+ return arith::ExtFOp::create(rewriter, loc, dstType, val,
+ arith::FastMathFlagsAttr{});
if (isa<IntegerType>(srcElementType) && isa<IntegerType>(dstElementType) &&
srcWidth < dstWidth)
diff --git a/mlir/lib/Dialect/Math/Transforms/ExtendToSupportedTypes.cpp b/mlir/lib/Dialect/Math/Transforms/ExtendToSupportedTypes.cpp
index bc262f84b26ac..550d8d9fa50c9 100644
--- a/mlir/lib/Dialect/Math/Transforms/ExtendToSupportedTypes.cpp
+++ b/mlir/lib/Dialect/Math/Transforms/ExtendToSupportedTypes.cpp
@@ -73,7 +73,9 @@ void mlir::math::populateExtendToSupportedTypesTypeConverter(
});
typeConverter.addTargetMaterialization(
[](OpBuilder &b, Type target, ValueRange input, Location loc) {
- auto extFOp = arith::ExtFOp::create(b, loc, target, input);
+ auto extFOp = arith::ExtFOp::create(b, loc, TypeRange{target},
+ ValueRange{input.front()},
+ arith::ExtFOp::Properties{});
extFOp.setFastmath(arith::FastMathFlags::contract);
return extFOp;
});
diff --git a/mlir/lib/Dialect/Math/Transforms/PolynomialApproximation.cpp b/mlir/lib/Dialect/Math/Transforms/PolynomialApproximation.cpp
index 9b367e0f84cb8..02fe79024ff5e 100644
--- a/mlir/lib/Dialect/Math/Transforms/PolynomialApproximation.cpp
+++ b/mlir/lib/Dialect/Math/Transforms/PolynomialApproximation.cpp
@@ -344,9 +344,12 @@ LogicalResult insertCasts(Operation *op, PatternRewriter &rewriter) {
Location loc = op->getLoc();
SmallVector<Value> operands;
for (auto operand : op->getOperands())
- operands.push_back(arith::ExtFOp::create(rewriter, loc, newType, operand));
- auto result =
- T::create(rewriter, loc, TypeRange{newType}, operands, op->getAttrs());
+ operands.push_back(arith::ExtFOp::create(rewriter, loc, TypeRange{newType},
+ ValueRange{operand},
+ arith::ExtFOp::Properties{}));
+ auto result = T::create(rewriter, loc, TypeRange{newType}, operands,
+ cast<T>(op).getProperties(),
+ op->getDiscardableAttrDictionary().getValue());
rewriter.replaceOpWithNewOp<arith::TruncFOp>(op, origType, result);
return success();
}
diff --git a/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp b/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp
index c6a2d1e97f644..4b5a4c1780507 100644
--- a/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp
+++ b/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp
@@ -2895,9 +2895,11 @@ void CollapseShapeOp::build(OpBuilder &b, OperationState &result, Value src,
auto srcType = llvm::cast<MemRefType>(src.getType());
MemRefType resultType =
CollapseShapeOp::computeCollapsedType(srcType, reassociation);
- result.addAttribute(::mlir::getReassociationAttrName(),
- getReassociationIndicesAttribute(b, reassociation));
- build(b, result, resultType, src, attrs);
+ buildPropertiesAndDiscardableAttributes(result, attrs);
+ result.getOrAddProperties<Properties>().reassociation =
+ getReassociationIndicesAttribute(b, reassociation);
+ result.addOperands(src);
+ result.addTypes(resultType);
}
LogicalResult CollapseShapeOp::verify() {
@@ -3830,8 +3832,10 @@ void TransposeOp::build(OpBuilder &b, OperationState &result, Value in,
// Compute result type.
MemRefType resultType = inferTransposeResultType(memRefType, permutationMap);
- result.addAttribute(TransposeOp::getPermutationAttrStrName(), permutation);
- build(b, result, resultType, in, attrs);
+ buildPropertiesAndDiscardableAttributes(result, attrs);
+ result.getOrAddProperties<Properties>().permutation = permutation;
+ result.addOperands(in);
+ result.addTypes(resultType);
}
// transpose $in $permutation attr-dict : type($in) `to` type(results)
diff --git a/mlir/lib/Dialect/MemRef/Transforms/MultiBuffer.cpp b/mlir/lib/Dialect/MemRef/Transforms/MultiBuffer.cpp
index 69adf0611b623..5fba3dfcf1fbd 100644
--- a/mlir/lib/Dialect/MemRef/Transforms/MultiBuffer.cpp
+++ b/mlir/lib/Dialect/MemRef/Transforms/MultiBuffer.cpp
@@ -224,6 +224,8 @@ mlir::memref::multiBuffer(RewriterBase &rewriter, memref::AllocOp allocOp,
Location loc = allocOp->getLoc();
OpBuilder::InsertionGuard g(rewriter);
rewriter.setInsertionPoint(allocOp);
+ memref::AllocOp::Properties properties = allocOp.getProperties();
+ properties.operandSegmentSizes = {0, 0};
auto mbAlloc = memref::AllocOp::create(
rewriter, loc, mbMemRefType, ValueRange{}, allocOp.getAlignmentAttr());
mbAlloc->setDiscardableAttrs(allocOp->getDiscardableAttrDictionary());
diff --git a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 360eac356cceb..f48ea70b9b260 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -224,7 +224,7 @@ struct MemRefPointerLikeModel
if (memrefTy.getRank() != 0)
return {};
- return memref::LoadOp::create(builder, loc, memrefValue);
+ return memref::LoadOp::create(builder, loc, memrefValue, ValueRange{});
}
bool genStore(Type pointer, OpBuilder &builder, Location loc,
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
index 9d8961da15743..6165c51c718bb 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -2772,11 +2772,11 @@ Value ACCCGToGPULowering::processPrivatize(acc::PrivatizeOp privatize) {
Value load;
if (threadYIsActive) {
Value threadYId = getThreadId(loc, gpu::Dimension::y);
- load = memref::LoadOp::create(rewriter, privatize->getLoc(), baseTy, alloca,
+ load = memref::LoadOp::create(rewriter, privatize->getLoc(), alloca,
ValueRange{threadYId});
} else {
- load =
- memref::LoadOp::create(rewriter, privatize->getLoc(), baseTy, alloca);
+ load = memref::LoadOp::create(rewriter, privatize->getLoc(), alloca,
+ ValueRange{});
}
rewriter.setInsertionPointAfter(load.getDefiningOp());
mapping.map(privatize.getResult(), load);
@@ -3944,7 +3944,8 @@ void ACCCGToGPULowering::processReductionCombineOp(acc::ReductionCombineOp op) {
// by the parent predicate_region processing.
// Reloading a grid-shared slot races with other blocks; record
// it and replace with the block-reduced register value in the fixup.
- auto srcLoad = memref::LoadOp::create(rewriter, loc, srcMemref);
+ auto srcLoad =
+ memref::LoadOp::create(rewriter, loc, srcMemref, ValueRange{});
pendingCombineReloads.push_back({srcMemref, srcLoad});
constructAtomicAccumulation(loc, destMemref, /*indices=*/{}, srcLoad,
kind);
@@ -3989,7 +3990,8 @@ void ACCCGToGPULowering::processCombineRegionOp(
return;
Value srcMemref = mapping.lookupOrDefault(accumulateOp.getMemref());
// Recorded and patched in the fixup to avoid the reload race.
- auto reductionLoad = memref::LoadOp::create(rewriter, loc, srcMemref);
+ auto reductionLoad =
+ memref::LoadOp::create(rewriter, loc, srcMemref, ValueRange{});
pendingCombineReloads.push_back({srcMemref, reductionLoad});
constructAtomicAccumulation(loc,
mapping.lookupOrDefault(op.getDestVar()),
@@ -4005,7 +4007,7 @@ void ACCCGToGPULowering::processCombineRegionOp(
if (isa<ComplexType>(memrefTy.getElementType())) {
Location loc = op.getLoc();
Value reductionResult =
- memref::LoadOp::create(rewriter, loc, privateMemref);
+ memref::LoadOp::create(rewriter, loc, privateMemref, ValueRange{});
arith::AtomicRMWKind kind = arith::AtomicRMWKind::addf;
op.getRegion().walk([&](Operation *innerOp) {
if (isa<complex::MulOp>(innerOp))
diff --git a/mlir/lib/Dialect/Shape/IR/Shape.cpp b/mlir/lib/Dialect/Shape/IR/Shape.cpp
index 4604786da7881..e3d49653ff629 100644
--- a/mlir/lib/Dialect/Shape/IR/Shape.cpp
+++ b/mlir/lib/Dialect/Shape/IR/Shape.cpp
@@ -593,8 +593,9 @@ struct RemoveDuplicateOperandsPattern : public OpRewritePattern<OpTy> {
// Reduce op to equivalent with unique operands.
if (unique.size() < op.getNumOperands()) {
- rewriter.replaceOpWithNewOp<OpTy>(op, op->getResultTypes(),
- unique.takeVector(), op->getAttrs());
+ rewriter.replaceOpWithNewOp<OpTy>(
+ op, op->getResultTypes(), unique.takeVector(), op.getProperties(),
+ op->getDiscardableAttrDictionary().getValue());
return success();
}
@@ -716,8 +717,9 @@ struct RemoveEmptyShapeOperandsPattern : public OpRewritePattern<OpTy> {
// Reduce op to equivalent without empty shape operands.
if (newOperands.size() < op.getNumOperands()) {
- rewriter.replaceOpWithNewOp<OpTy>(op, op->getResultTypes(), newOperands,
- op->getAttrs());
+ rewriter.replaceOpWithNewOp<OpTy>(
+ op, op->getResultTypes(), newOperands, op.getProperties(),
+ op->getDiscardableAttrDictionary().getValue());
return success();
}
@@ -786,8 +788,9 @@ struct BroadcastFoldConstantOperandsPattern
newShapeOperands.push_back(
ConstShapeOp::create(rewriter, op.getLoc(), foldedConstantOperandsTy,
rewriter.getIndexTensorAttr(foldedConstantShape)));
- rewriter.replaceOpWithNewOp<BroadcastOp>(op, op.getType(),
- newShapeOperands);
+ rewriter.replaceOpWithNewOp<BroadcastOp>(
+ op, TypeRange{op.getType()}, newShapeOperands, op.getProperties(),
+ op->getDiscardableAttrDictionary().getValue());
return success();
}
};
@@ -819,7 +822,9 @@ struct CanonicalizeCastExtentTensorOperandsPattern
// Rewrite op if any change required.
if (!anyChange)
return failure();
- rewriter.replaceOpWithNewOp<OpTy>(op, op->getResultTypes(), newOperands);
+ rewriter.replaceOpWithNewOp<OpTy>(
+ op, op->getResultTypes(), newOperands, op.getProperties(),
+ op->getDiscardableAttrDictionary().getValue());
return success();
}
};
@@ -850,7 +855,7 @@ struct BroadcastConcretizeResultTypePattern
auto newOp = BroadcastOp::create(rewriter, op.getLoc(),
getExtentTensorType(getContext(), maxRank),
- op.getShapes());
+ op.getShapes(), /*error=*/nullptr);
rewriter.replaceOpWithNewOp<tensor::CastOp>(op, op.getType(), newOp);
return success();
}
@@ -1343,7 +1348,8 @@ std::optional<int64_t> GetExtentOp::getConstantDim() {
}
OpFoldResult GetExtentOp::fold(FoldAdaptor adaptor) {
- auto elements = llvm::dyn_cast_if_present<DenseIntElementsAttr>(adaptor.getShape());
+ auto elements =
+ llvm::dyn_cast_if_present<DenseIntElementsAttr>(adaptor.getShape());
if (!elements)
return nullptr;
std::optional<int64_t> dim = getConstantDim();
@@ -1485,7 +1491,8 @@ bool mlir::shape::MeetOp::isCompatibleReturnTypes(TypeRange l, TypeRange r) {
//===----------------------------------------------------------------------===//
OpFoldResult shape::RankOp::fold(FoldAdaptor adaptor) {
- auto shape = llvm::dyn_cast_if_present<DenseIntElementsAttr>(adaptor.getShape());
+ auto shape =
+ llvm::dyn_cast_if_present<DenseIntElementsAttr>(adaptor.getShape());
if (!shape)
return {};
int64_t rank = shape.getNumElements();
@@ -1727,12 +1734,14 @@ struct ShapeOfOpToConstShapeOp : public OpRewritePattern<shape::ShapeOfOp> {
// Canonicalize
//
-// %0 = tensor.reshape %input(%shape) : (tensor<*xf32>, tensor<?xindex>) -> tensor<*xf32>
+// %0 = tensor.reshape %input(%shape) : (tensor<*xf32>, tensor<?xindex>) ->
+// tensor<*xf32>
// %1 = shape.shape_of %0 : tensor<*xf32> -> tensor<?xindex>
//
// to
//
-// %0 = tensor.reshape %input(%shape) : (tensor<*xf32>, tensor<?xindex>) -> tensor<*xf32>
+// %0 = tensor.reshape %input(%shape) : (tensor<*xf32>, tensor<?xindex>) ->
+// tensor<*xf32>
// %1 = %shape
//
struct ShapeOfFromReshape : public OpRewritePattern<shape::ShapeOfOp> {
@@ -1906,8 +1915,9 @@ LogicalResult SplitAtOp::fold(FoldAdaptor adaptor,
SmallVectorImpl<OpFoldResult> &results) {
if (!adaptor.getOperand() || !adaptor.getIndex())
return failure();
- auto shapeVec = llvm::to_vector<6>(
- llvm::cast<DenseIntElementsAttr>(adaptor.getOperand()).getValues<int64_t>());
+ auto shapeVec =
+ llvm::to_vector<6>(llvm::cast<DenseIntElementsAttr>(adaptor.getOperand())
+ .getValues<int64_t>());
auto shape = llvm::ArrayRef(shapeVec);
auto splitPoint = llvm::cast<IntegerAttr>(adaptor.getIndex()).getInt();
// Verify that the split point is in the correct range.
@@ -1931,8 +1941,9 @@ OpFoldResult ToExtentTensorOp::fold(FoldAdaptor adaptor) {
if (!adaptor.getInput())
return OpFoldResult();
Builder builder(getContext());
- auto shape = llvm::to_vector<6>(
- llvm::cast<DenseIntElementsAttr>(adaptor.getInput()).getValues<int64_t>());
+ auto shape =
+ llvm::to_vector<6>(llvm::cast<DenseIntElementsAttr>(adaptor.getInput())
+ .getValues<int64_t>());
auto type = RankedTensorType::get({static_cast<int64_t>(shape.size())},
builder.getIndexType());
return DenseIntElementsAttr::get(type, shape);
diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseVectorization.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseVectorization.cpp
index 548eef578c4ff..b5ec17c88b205 100644
--- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseVectorization.cpp
+++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseVectorization.cpp
@@ -385,9 +385,11 @@ static bool vectorizeSubscripts(PatternRewriter &rewriter, scf::ForOp forOp,
}
#define UNAOP(xxx) \
- if (isa<xxx>(def)) { \
- if (codegen) \
- vexp = xxx::create(rewriter, loc, vx); \
+ if (auto x = dyn_cast<xxx>(def)) { \
+ if (codegen) { \
+ vexp = xxx::create(rewriter, loc, ValueRange{vx}, x.getProperties(), \
+ x->getDiscardableAttrDictionary().getValue()); \
+ } \
return true; \
}
@@ -395,15 +397,19 @@ static bool vectorizeSubscripts(PatternRewriter &rewriter, scf::ForOp forOp,
if (auto x = dyn_cast<xxx>(def)) { \
if (codegen) { \
VectorType vtp = vectorType(vl, x.getType()); \
- vexp = xxx::create(rewriter, loc, vtp, vx); \
+ vexp = xxx::create(rewriter, loc, TypeRange{vtp}, ValueRange{vx}, \
+ x.getProperties(), \
+ x->getDiscardableAttrDictionary().getValue()); \
} \
return true; \
}
#define BINOP(xxx) \
- if (isa<xxx>(def)) { \
- if (codegen) \
- vexp = xxx::create(rewriter, loc, vx, vy); \
+ if (auto x = dyn_cast<xxx>(def)) { \
+ if (codegen) { \
+ vexp = xxx::create(rewriter, loc, ValueRange{vx, vy}, x.getProperties(), \
+ x->getDiscardableAttrDictionary().getValue()); \
+ } \
return true; \
}
diff --git a/mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp b/mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp
index 5847fecc45404..f388a2e44ebdb 100644
--- a/mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp
+++ b/mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp
@@ -1674,7 +1674,8 @@ Value Merger::buildExp(RewriterBase &rewriter, Location loc, ExprId e, Value v0,
case TensorExp::Kind::kTruncF:
return arith::TruncFOp::create(rewriter, loc, inferType(e, v0), v0);
case TensorExp::Kind::kExtF:
- return arith::ExtFOp::create(rewriter, loc, inferType(e, v0), v0);
+ return arith::ExtFOp::create(rewriter, loc, TypeRange{inferType(e, v0)},
+ ValueRange{v0}, arith::ExtFOp::Properties{});
case TensorExp::Kind::kCastFS:
return arith::FPToSIOp::create(rewriter, loc, inferType(e, v0), v0);
case TensorExp::Kind::kCastFU:
diff --git a/mlir/lib/Dialect/Tensor/IR/TensorOps.cpp b/mlir/lib/Dialect/Tensor/IR/TensorOps.cpp
index fb14c1953fd22..5010cfb04589c 100644
--- a/mlir/lib/Dialect/Tensor/IR/TensorOps.cpp
+++ b/mlir/lib/Dialect/Tensor/IR/TensorOps.cpp
@@ -2055,9 +2055,11 @@ void CollapseShapeOp::build(OpBuilder &b, OperationState &result, Value src,
auto resultType =
RankedTensorType::get(collapsedType.getShape(), srcType.getElementType(),
srcType.getEncoding());
- result.addAttribute(getReassociationAttrStrName(),
- getReassociationIndicesAttribute(b, reassociation));
- build(b, result, resultType, src, attrs);
+ buildPropertiesAndDiscardableAttributes(result, attrs);
+ result.getOrAddProperties<Properties>().reassociation =
+ getReassociationIndicesAttribute(b, reassociation);
+ result.addOperands(src);
+ result.addTypes(resultType);
}
template <typename TensorReshapeOp, bool isExpansion = std::is_same<
diff --git a/mlir/lib/Dialect/Tosa/Transforms/TosaNarrowTypes.cpp b/mlir/lib/Dialect/Tosa/Transforms/TosaNarrowTypes.cpp
index 1ec465124dfb4..2968c2bc7ae44 100644
--- a/mlir/lib/Dialect/Tosa/Transforms/TosaNarrowTypes.cpp
+++ b/mlir/lib/Dialect/Tosa/Transforms/TosaNarrowTypes.cpp
@@ -504,7 +504,8 @@ class ConvertCastOpWithBoundsChecking
return failure();
rewriter.replaceOpWithNewOp<tosa::CastOp>(
- op, typeConverter->convertType(resultType), adaptor.getInput(),
+ op, TypeRange{typeConverter->convertType(resultType)},
+ ValueRange{adaptor.getInput()}, op.getProperties(),
op->getDiscardableAttrDictionary().getValue());
return success();
}
diff --git a/mlir/lib/Dialect/Vector/Transforms/LowerVectorContract.cpp b/mlir/lib/Dialect/Vector/Transforms/LowerVectorContract.cpp
index f6fff0ffe323a..1530367a349aa 100644
--- a/mlir/lib/Dialect/Vector/Transforms/LowerVectorContract.cpp
+++ b/mlir/lib/Dialect/Vector/Transforms/LowerVectorContract.cpp
@@ -393,7 +393,8 @@ struct UnrolledOuterProductGenerator
if (vecType)
promotedType = vecType.clone(promotedType);
if (isa<FloatType>(dstElementType))
- return arith::ExtFOp::create(rewriter, loc, promotedType, v);
+ return arith::ExtFOp::create(rewriter, loc, promotedType, v,
+ arith::FastMathFlagsAttr{});
return arith::ExtSIOp::create(rewriter, loc, promotedType, v);
}
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp
index 19d02e02b02ba..7d97fbcb8e2af 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp
@@ -1754,9 +1754,10 @@ ResolveLayoutConflicts::resolveTensorDescConsumer(OpOperand &operand) {
currTDescType.getElementType(), currTDescType.getEncoding(),
expectedLayout);
xegpu::CreateNdDescOp newOp = xegpu::CreateNdDescOp::create(
- builder, consumerOp->getLoc(), newTensorDescType,
+ builder, consumerOp->getLoc(), TypeRange{newTensorDescType},
conflictingCreateNdOp->getOperands(),
- conflictingCreateNdOp->getAttrs());
+ conflictingCreateNdOp.getProperties(),
+ conflictingCreateNdOp->getDiscardableAttrDictionary().getValue());
// Replace the tensor descriptor operand in the consumer op with the new
// tensor descriptor.
consumerOp->replaceUsesOfWith(tdescValue, newOp.getResult());
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp
index bc85140f9f121..9b6ec9b7b2a06 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp
@@ -128,8 +128,9 @@ struct SgToLaneCreateNdDesc
return failure();
auto newOp = xegpu::CreateNdDescOp::create(
- rewriter, op.getLoc(), resultType.dropLayouts(), op.getOperands(),
- op->getAttrs());
+ rewriter, op.getLoc(), TypeRange{resultType.dropLayouts()},
+ op.getOperands(), op.getProperties(),
+ op->getDiscardableAttrDictionary().getValue());
rewriter.replaceOp(op, newOp.getResult());
return success();
}
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp
index 5914e23f1f11a..b898d87481635 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp
@@ -43,6 +43,11 @@ unrollByTile(SmallVector<OpFoldResult> mixedOffsets,
const std::function<Value(SmallVector<OpFoldResult>)> &createOp,
Location loc, PatternRewriter &rewriter);
+static xegpu::DistributeLayoutAttr
+dropInstData(xegpu::DistributeLayoutAttr layout) {
+ return layout ? layout.dropInstData() : nullptr;
+}
+
template <typename SourceOp>
struct UnrollPattern : public OpRewritePattern<SourceOp> {
UnrollPattern(MLIRContext *context, const xegpu::UnrollOptions &options,
@@ -428,9 +433,14 @@ struct UnrollDpasOp : public UnrollPattern<xegpu::DpasOp> {
if (tmpC)
operands.push_back(tmpC);
+ auto properties = op.getProperties();
+ properties.layout_a = dropInstData(properties.layout_a);
+ properties.layout_b = dropInstData(properties.layout_b);
+ properties.layout_cd = dropInstData(properties.layout_cd);
tmpC = xegpu::DpasOp::create(
- rewriter, loc, vecTy, operands,
- xegpu::dropInstDataOnAttrs(op->getAttrs()));
+ rewriter, loc, TypeRange{vecTy}, operands, properties,
+ xegpu::dropInstDataOnAttrs(
+ op->getDiscardableAttrDictionary().getValue()));
}
newOps.push_back(tmpC);
}
@@ -530,9 +540,16 @@ struct UnrollDpasMxOp : public UnrollPattern<xegpu::DpasMxOp> {
operands.push_back(
bScaleVals[batch * (kIters * nIters) + k * nIters + j]);
+ auto properties = op.getProperties();
+ properties.layout_a = dropInstData(properties.layout_a);
+ properties.layout_b = dropInstData(properties.layout_b);
+ properties.layout_cd = dropInstData(properties.layout_cd);
+ properties.layout_a_scale = dropInstData(properties.layout_a_scale);
+ properties.layout_b_scale = dropInstData(properties.layout_b_scale);
newDpasMxOp = xegpu::DpasMxOp::create(
- rewriter, loc, vecTy, operands,
- xegpu::dropInstDataOnAttrs(op->getAttrs()));
+ rewriter, loc, TypeRange{vecTy}, operands, properties,
+ xegpu::dropInstDataOnAttrs(
+ op->getDiscardableAttrDictionary().getValue()));
tmpC = newDpasMxOp.getResult();
}
newOps.push_back(newDpasMxOp);
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
index 0e8a386fb08b6..18ef34170d9fe 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
@@ -339,7 +339,9 @@ struct WgToSgDpasOp : public OpConversionPattern<xegpu::DpasOp> {
resShape.push_back(aVecShape[aVecShape.size() - 2]);
resShape.push_back(bVecShape[bVecShape.size() - 1]);
VectorType resTy = VectorType::get(resShape, resultTy.getElementType());
- auto newDpasOp = xegpu::DpasOp::create(rewriter, loc, resTy, operands);
+ auto newDpasOp = xegpu::DpasOp::create(
+ rewriter, loc, resTy, operands[0], operands[1], tmpC,
+ /*layout_a=*/nullptr, /*layout_b=*/nullptr, /*layout_cd=*/nullptr);
newDpasOp.setLayoutCdAttr(layoutCd.dropSgLayoutAndData());
newDpasOp.setLayoutAAttr(layoutA.dropSgLayoutAndData());
newDpasOp.setLayoutBAttr(layoutB.dropSgLayoutAndData());
diff --git a/mlir/lib/TableGen/Operator.cpp b/mlir/lib/TableGen/Operator.cpp
index 148c3408b707d..767f00db2fe81 100644
--- a/mlir/lib/TableGen/Operator.cpp
+++ b/mlir/lib/TableGen/Operator.cpp
@@ -274,6 +274,35 @@ const Trait *Operator::getTrait(StringRef trait) const {
return nullptr;
}
+bool Operator::hasNonEmptyProperties() const {
+ if (!properties.empty())
+ return true;
+ if (getTrait("::mlir::OpTrait::AttrSizedOperandSegments") ||
+ getTrait("::mlir::OpTrait::AttrSizedResultSegments"))
+ return true;
+ return llvm::any_of(attributes, [](const NamedAttribute &attr) {
+ return !attr.attr.isDerivedAttr();
+ });
+}
+
+SmallVector<StringRef> Operator::getInherentAttrNames() const {
+ SmallVector<StringRef> names;
+ for (const NamedAttribute &attr : attributes)
+ if (!attr.attr.isDerivedAttr())
+ names.push_back(attr.name);
+ for (const NamedProperty &property : properties)
+ names.push_back(property.name);
+ if (getTrait("::mlir::OpTrait::AttrSizedOperandSegments")) {
+ names.push_back(operandSegmentAttrName);
+ names.push_back(legacyOperandSegmentAttrName);
+ }
+ if (getTrait("::mlir::OpTrait::AttrSizedResultSegments")) {
+ names.push_back(resultSegmentAttrName);
+ names.push_back(legacyResultSegmentAttrName);
+ }
+ return names;
+}
+
auto Operator::region_begin() const -> const_region_iterator {
return regions.begin();
}
diff --git a/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp b/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp
index 66c77c61597d5..acb43b6e04bd6 100644
--- a/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp
+++ b/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp
@@ -747,8 +747,17 @@ Deserializer::processOp<spirv::CopyMemoryOp>(ArrayRef<uint32_t> words) {
}
Location loc = createFileLineColLoc(opBuilder);
- spirv::CopyMemoryOp::create(opBuilder, loc, resultTypes, operands,
- attributes);
+ spirv::CopyMemoryOp::Properties properties{};
+ spirv::CopyMemoryOp::populateDefaultProperties(
+ OperationName(spirv::CopyMemoryOp::getOperationName(),
+ opBuilder.getContext()),
+ properties);
+ if (failed(spirv::CopyMemoryOp::setPropertiesFromAttr(
+ properties, opBuilder.getDictionaryAttr(attributes),
+ [&]() { return emitError(loc); })))
+ return failure();
+ spirv::CopyMemoryOp::create(opBuilder, loc, resultTypes, operands, properties,
+ /*discardableAttributes=*/{});
return success();
}
diff --git a/mlir/lib/Tools/PDLL/CodeGen/MLIRGen.cpp b/mlir/lib/Tools/PDLL/CodeGen/MLIRGen.cpp
index 08c085b5f279c..6934258647f91 100644
--- a/mlir/lib/Tools/PDLL/CodeGen/MLIRGen.cpp
+++ b/mlir/lib/Tools/PDLL/CodeGen/MLIRGen.cpp
@@ -441,7 +441,8 @@ Value CodeGen::genExprImpl(const ast::MemberAccessExpr *expr) {
if (isa<pdl::ValueType>(mlirType))
return pdl::ResultOp::create(builder, loc, mlirType, parentExprs[0],
builder.getI32IntegerAttr(0));
- return pdl::ResultsOp::create(builder, loc, mlirType, parentExprs[0]);
+ return pdl::ResultsOp::create(builder, loc, mlirType, parentExprs[0],
+ /*index=*/nullptr);
}
const ods::Operation *odsOp = opType.getODSOperation();
diff --git a/mlir/test/Dialect/SPIRV/IR/target-env.mlir b/mlir/test/Dialect/SPIRV/IR/target-env.mlir
index 7d807bbbfbd4f..48421076402a9 100644
--- a/mlir/test/Dialect/SPIRV/IR/target-env.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/target-env.mlir
@@ -140,7 +140,8 @@ func.func @sdot_scalar_i32_i32_capabilities(%operand: i32) -> i32 attributes {
[DotProduct, DotProductInput4x8BitPacked], [SPV_KHR_integer_dot_product]>, #spirv.resource_limits<>>
} {
// CHECK: spirv.SDot
- %0 = "test.convert_to_sdot_op"(%operand, %operand) {format = #spirv.packed_vector_format<PackedVectorFormat4x8Bit>}: (i32, i32) -> (i32)
+ // CHECK-SAME: test.marker = "keep"
+ %0 = "test.convert_to_sdot_op"(%operand, %operand) {format = #spirv.packed_vector_format<PackedVectorFormat4x8Bit>, test.marker = "keep"}: (i32, i32) -> (i32)
return %0: i32
}
diff --git a/mlir/test/Dialect/Tosa/tosa-narrow-i64-to-i32.mlir b/mlir/test/Dialect/Tosa/tosa-narrow-i64-to-i32.mlir
index 4ff8a38548c81..a8b5de834de8c 100644
--- a/mlir/test/Dialect/Tosa/tosa-narrow-i64-to-i32.mlir
+++ b/mlir/test/Dialect/Tosa/tosa-narrow-i64-to-i32.mlir
@@ -31,8 +31,8 @@ func.func @test_i64_const() -> tensor<2xi64> {
func.func @test_i64_argmax_cast(%arg0: tensor<1x513x513x19xi8>) -> tensor<1x513x513xf32> {
// COMMON: %[[ARGMAX:.*]] = tosa.argmax %arg0 {axis = 3 : i32} : (tensor<1x513x513x19xi8>) -> tensor<1x513x513xi32>
%0 = tosa.argmax %arg0 {axis = 3 : i32} : (tensor<1x513x513x19xi8>) -> tensor<1x513x513xi64>
- // COMMON: tosa.cast %[[ARGMAX]] : (tensor<1x513x513xi32>) -> tensor<1x513x513xf32>
- %1 = tosa.cast %0 {input_unsigned = false} : (tensor<1x513x513xi64>) -> tensor<1x513x513xf32>
+ // COMMON: tosa.cast %[[ARGMAX]] {test.marker = "keep"} : (tensor<1x513x513xi32>) -> tensor<1x513x513xf32>
+ %1 = tosa.cast %0 {input_unsigned = false, test.marker = "keep"} : (tensor<1x513x513xi64>) -> tensor<1x513x513xf32>
return %1 : tensor<1x513x513xf32>
}
diff --git a/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp b/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp
index 8cdb9041d196a..55e3581a3a58b 100644
--- a/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp
+++ b/mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp
@@ -208,9 +208,24 @@ struct ConvertToIntegerDotProd : RewritePattern {
LogicalResult matchAndRewrite(Operation *op,
PatternRewriter &rewriter) const override {
- rewriter.replaceOpWithNewOp<SPIRVOp>(
- op, op->getResultTypes(), op->getOperands(),
- op->getDiscardableAttrDictionary().getValue());
+ typename SPIRVOp::Properties properties{};
+ SPIRVOp::populateDefaultProperties(
+ OperationName(SPIRVOp::getOperationName(), rewriter.getContext()),
+ properties);
+ if (failed(SPIRVOp::setPropertiesFromAttr(
+ properties, op->getDiscardableAttrDictionary(),
+ [&]() { return op->emitError("invalid SPIR-V properties"); })))
+ return failure();
+ auto propertiesAttr = dyn_cast_or_null<DictionaryAttr>(
+ SPIRVOp::getPropertiesAsAttr(rewriter.getContext(), properties));
+ SmallVector<NamedAttribute> discardableAttrs;
+ for (NamedAttribute attr : op->getDiscardableAttrDictionary().getValue()) {
+ if (!propertiesAttr || !propertiesAttr.contains(attr.getName()))
+ discardableAttrs.push_back(attr);
+ }
+ rewriter.replaceOpWithNewOp<SPIRVOp>(op, op->getResultTypes(),
+ op->getOperands(), properties,
+ discardableAttrs);
return success();
}
};
diff --git a/mlir/test/lib/Dialect/Test/TestPatterns.cpp b/mlir/test/lib/Dialect/Test/TestPatterns.cpp
index fbc5ab76d9ba1..52d8f4189fca1 100644
--- a/mlir/test/lib/Dialect/Test/TestPatterns.cpp
+++ b/mlir/test/lib/Dialect/Test/TestPatterns.cpp
@@ -112,7 +112,8 @@ struct FoldingPattern : public RewritePattern {
// should not crash when attempting to recover the (unchanged) operation
// result.
Value result = rewriter.createOrFold<TestOpInPlaceFold>(
- op->getLoc(), rewriter.getIntegerType(32), op->getOperand(0));
+ op->getLoc(), rewriter.getIntegerType(32), op->getOperand(0),
+ IntegerAttr{});
assert(result);
rewriter.replaceOp(op, result);
return success();
diff --git a/mlir/test/mlir-tblgen/op-decl-and-defs.td b/mlir/test/mlir-tblgen/op-decl-and-defs.td
index 28716cf229577..32c4bb88d188b 100644
--- a/mlir/test/mlir-tblgen/op-decl-and-defs.td
+++ b/mlir/test/mlir-tblgen/op-decl-and-defs.td
@@ -135,8 +135,11 @@ def NS_AOp : NS_Op<"a_op", [IsolatedFromAbove, IsolatedFromAbove]> {
// CHECK: static void build(::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::Value a, ::mlir::ValueRange b, uint32_t attr1, /*optional*/::mlir::FloatAttr some_attr2, unsigned someRegionsCount);
// CHECK: static AOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::Value a, ::mlir::ValueRange b, uint32_t attr1, /*optional*/::mlir::FloatAttr some_attr2, unsigned someRegionsCount);
// CHECK: static AOp create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::TypeRange resultTypes, ::mlir::Value a, ::mlir::ValueRange b, uint32_t attr1, /*optional*/::mlir::FloatAttr some_attr2, unsigned someRegionsCount);
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes, unsigned numRegions)
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static AOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes, unsigned numRegions)
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static AOp create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes, unsigned numRegions)
// CHECK-NEXT: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes, unsigned numRegions)
// CHECK-NEXT: static AOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes, unsigned numRegions)
@@ -266,14 +269,20 @@ def NS_FirstAttrDerivedOp : NS_Op<"first_attr_derived",
}
// CHECK-LABEL: class FirstAttrDerivedOp :
-// CHECK: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
+// CHECK-NEXT: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static FirstAttrDerivedOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static FirstAttrDerivedOp create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
// CHECK-NEXT: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes = {});
// CHECK-NEXT: static FirstAttrDerivedOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes = {});
// CHECK-NEXT: static FirstAttrDerivedOp create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static void build(::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &odsState, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static FirstAttrDerivedOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static FirstAttrDerivedOp create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
// CHECK-NEXT: static void build(::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &odsState, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes = {});
// CHECK-NEXT: static FirstAttrDerivedOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes = {});
@@ -320,6 +329,7 @@ def NS_HCollectiveParamsSuppress0Op : NS_Op<"op_collective_suppress0", []> {
// CHECK-NOT: static void build(::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &odsState, ::mlir::TypeRange b, ::mlir::ValueRange a);
// CHECK-NOT: static HCollectiveParamsSuppress0Op create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange b, ::mlir::ValueRange a);
// CHECK-NOT: static HCollectiveParamsSuppress0Op create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::TypeRange b, ::mlir::ValueRange a);
+// CHECK-NOT: use the overload taking Properties and discardableAttributes instead
// CHECK: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
// CHECK: static HCollectiveParamsSuppress0Op create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
// CHECK: static HCollectiveParamsSuppress0Op create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
@@ -479,11 +489,17 @@ def NS_LOp : NS_Op<"op_with_same_operands_and_result_types_unwrapped_attr", [Sam
// CHECK: static LOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::Value a, ::mlir::Value b, uint32_t attr1);
// CHECK: static LOp create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::TypeRange resultTypes, ::mlir::Value a, ::mlir::Value b, uint32_t attr1);
-// CHECK: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
+// CHECK-NEXT: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static LOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static LOp create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static void build(::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &odsState, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static LOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static LOp create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
// CHECK-NEXT: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes = {});
// CHECK-NEXT: static LOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes = {});
@@ -583,11 +599,17 @@ def _TypeInferredPropOp : NS_Op<"type_inferred_prop_op_with_properties", [
let hasCustomAssemblyFormat = 1;
}
-// CHECK: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
+// CHECK-NEXT: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static _TypeInferredPropOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static _TypeInferredPropOp create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static void build(::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &odsState, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static _TypeInferredPropOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
+// CHECK{LITERAL}: [[deprecated("use the overload taking Properties and discardableAttributes instead")]]
// CHECK-NEXT: static _TypeInferredPropOp create(::mlir::ImplicitLocOpBuilder &builder, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes = {});
// CHECK-NEXT: static void build(::mlir::OpBuilder &, ::mlir::OperationState &odsState, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes = {});
// CHECK-NEXT: static _TypeInferredPropOp create(::mlir::OpBuilder &builder, ::mlir::Location location, ::mlir::TypeRange resultTypes, ::mlir::ValueRange operands, const Properties &properties, ::llvm::ArrayRef<::mlir::NamedAttribute> discardableAttributes = {});
diff --git a/mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp b/mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp
index 84048f76c630c..da3446b6349a1 100644
--- a/mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp
+++ b/mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp
@@ -62,20 +62,13 @@ static const char *const odsBuilder = "odsBuilder";
static const char *const builderOpState = "odsState";
static const char *const builderOpStateProperties =
"odsState.getOrAddProperties<Properties>()";
+static constexpr StringLiteral legacyBuilderDeprecation =
+ "use the overload taking Properties and discardableAttributes instead";
static const char *const propertyStorage = "propStorage";
static const char *const propertyValue = "propValue";
static const char *const propertyAttr = "propAttr";
static const char *const propertyDiag = "emitError";
-/// The names of the implicit attributes that contain variadic operand and
-/// result segment sizes.
-static const char *const operandSegmentAttrName = "operandSegmentSizes";
-static const char *const resultSegmentAttrName = "resultSegmentSizes";
-static constexpr StringLiteral legacyOperandSegmentAttrName =
- "operand_segment_sizes";
-static constexpr StringLiteral legacyResultSegmentAttrName =
- "result_segment_sizes";
-
/// Code for an Op to lookup an attribute. Uses cached identifiers and subrange
/// lookup.
///
@@ -431,18 +424,7 @@ class OpOrAdaptorHelper {
/// Returns whether the operation will have a non-empty `Properties` struct.
bool hasNonEmptyPropertiesStruct() const {
- if (!op.getProperties().empty())
- return true;
- if (!hasProperties())
- return false;
- if (op.getTrait("::mlir::OpTrait::AttrSizedOperandSegments") ||
- op.getTrait("::mlir::OpTrait::AttrSizedResultSegments"))
- return true;
- return llvm::any_of(getAttrMetadata(),
- [](const std::pair<StringRef, AttributeMetadata> &it) {
- return !it.second.constraint ||
- !it.second.constraint->isDerivedAttr();
- });
+ return hasProperties() && op.hasNonEmptyProperties();
}
std::optional<NamedProperty> &getOperandSegmentsSize() {
@@ -666,7 +648,8 @@ class OpEmitter {
// Generates the build() method that takes each operand/attribute
// as a stand-alone parameter.
void genSeparateArgParamBuilder();
- void genInlineCreateBody(const SmallVector<MethodParameter> ¶mList);
+ void genInlineCreateBody(const SmallVector<MethodParameter> ¶mList,
+ bool deprecated = false);
// Generates the build() method that takes each operand/attribute as a
// stand-alone parameter. The generated build() method uses first operand's
@@ -2548,7 +2531,7 @@ static bool canInferType(const Operator &op) {
}
void OpEmitter::genInlineCreateBody(
- const SmallVector<MethodParameter> ¶mList) {
+ const SmallVector<MethodParameter> ¶mList, bool deprecated) {
SmallVector<MethodParameter> createParamListOpBuilder;
SmallVector<MethodParameter> createParamListImplicitLocOpBuilder;
SmallVector<llvm::StringRef, 4> nonBuilderStateArgsList;
@@ -2579,6 +2562,12 @@ void OpEmitter::genInlineCreateBody(
createParamListOpBuilder);
auto *cImplicitLoc = opClass.addStaticMethod(
opClass.getClassName(), "create", createParamListImplicitLocOpBuilder);
+ if (deprecated) {
+ if (cWithLoc)
+ cWithLoc->setDeprecated(legacyBuilderDeprecation);
+ if (cImplicitLoc)
+ cImplicitLoc->setDeprecated(legacyBuilderDeprecation);
+ }
std::string nonBuilderStateArgs = "";
if (!nonBuilderStateArgsList.empty()) {
llvm::raw_string_ostream nonBuilderStateArgsOS(nonBuilderStateArgs);
@@ -2718,21 +2707,7 @@ void OpEmitter::genLegacyPropertiesBuilderHelper() {
if (!emitHelper.hasNonEmptyPropertiesStruct())
return;
- SmallVector<StringRef> inherentNames;
- for (const ConstArgument &attrOrProperty : getAttrOrProperties()) {
- if (const auto *namedAttr =
- dyn_cast_if_present<const AttributeMetadata *>(attrOrProperty))
- inherentNames.push_back(namedAttr->attrName);
- else
- inherentNames.push_back(
- cast<const NamedProperty *>(attrOrProperty)->name);
- }
- if (emitHelper.getOperandSegmentsSize()) {
- inherentNames.push_back(legacyOperandSegmentAttrName);
- }
- if (emitHelper.getResultSegmentsSize()) {
- inherentNames.push_back(legacyResultSegmentAttrName);
- }
+ SmallVector<StringRef> inherentNames = op.getInherentAttrNames();
llvm::sort(inherentNames);
inherentNames.erase(llvm::unique(inherentNames), inherentNames.end());
@@ -2815,7 +2790,11 @@ void OpEmitter::genUseOperandAsResultTypeCollectiveParamBuilder(
// If the builder is redundant, skip generating the method
if (!m)
return;
- genInlineCreateBody(paramList);
+ bool deprecated = kind == CollectiveBuilderKind::AttrDict &&
+ emitHelper.hasNonEmptyPropertiesStruct();
+ if (deprecated)
+ m->setDeprecated(legacyBuilderDeprecation);
+ genInlineCreateBody(paramList, deprecated);
auto &body = m->body();
// Operands
@@ -2925,7 +2904,11 @@ void OpEmitter::genInferredTypeCollectiveParamBuilder(
// If the builder is redundant, skip generating the method
if (!m)
return;
- genInlineCreateBody(paramList);
+ bool deprecated = kind == CollectiveBuilderKind::AttrDict &&
+ emitHelper.hasNonEmptyPropertiesStruct();
+ if (deprecated)
+ m->setDeprecated(legacyBuilderDeprecation);
+ genInlineCreateBody(paramList, deprecated);
auto &body = m->body();
int numResults = op.getNumResults();
@@ -3029,7 +3012,11 @@ void OpEmitter::genUseAttrAsResultTypeCollectiveParamBuilder(
// If the builder is redundant, skip generating the method
if (!m)
return;
- genInlineCreateBody(paramList);
+ bool deprecated = kind == CollectiveBuilderKind::AttrDict &&
+ emitHelper.hasNonEmptyPropertiesStruct();
+ if (deprecated)
+ m->setDeprecated(legacyBuilderDeprecation);
+ genInlineCreateBody(paramList, deprecated);
auto &body = m->body();
@@ -3197,7 +3184,11 @@ void OpEmitter::genCollectiveParamBuilder(CollectiveBuilderKind kind) {
// If the builder is redundant, skip generating the method
if (!m)
return;
- genInlineCreateBody(paramList);
+ bool deprecated = kind == CollectiveBuilderKind::AttrDict &&
+ emitHelper.hasNonEmptyPropertiesStruct();
+ if (deprecated)
+ m->setDeprecated(legacyBuilderDeprecation);
+ genInlineCreateBody(paramList, deprecated);
auto &body = m->body();
// Operands
diff --git a/mlir/tools/mlir-tblgen/SPIRVUtilsGen.cpp b/mlir/tools/mlir-tblgen/SPIRVUtilsGen.cpp
index f9f25ebb1e663..d42aef47ba066 100644
--- a/mlir/tools/mlir-tblgen/SPIRVUtilsGen.cpp
+++ b/mlir/tools/mlir-tblgen/SPIRVUtilsGen.cpp
@@ -1066,10 +1066,42 @@ static void emitDeserializationFunction(const Record *attrClass,
emitDecorationDeserialization(op, " ", valueID, attributes, os);
os << formatv(" Location loc = createFileLineColLoc(opBuilder);\n");
- os << formatv(" auto {1} = {0}::create(opBuilder, loc, {2}, {3}, {4}); "
- "(void){1};\n",
- op.getQualCppClassName(), opVar, resultTypes, operands,
- attributes);
+ if (!op.hasNonEmptyProperties()) {
+ os << formatv(" auto {1} = {0}::create(opBuilder, loc, {2}, {3}, {4}); "
+ "(void){1};\n",
+ op.getQualCppClassName(), opVar, resultTypes, operands,
+ attributes);
+ } else {
+ os << formatv(" {0}::Properties properties{{};\n",
+ op.getQualCppClassName());
+ os << formatv(" {0}::populateDefaultProperties(\n"
+ " OperationName({0}::getOperationName(), "
+ "opBuilder.getContext()), properties);\n",
+ op.getQualCppClassName());
+ os << formatv(" if (failed({0}::setPropertiesFromAttr(\n"
+ " properties, opBuilder.getDictionaryAttr({1}),\n"
+ " [&]() {{ return emitError(loc); })))\n"
+ " return failure();\n",
+ op.getQualCppClassName(), attributes);
+ os << " SmallVector<NamedAttribute> discardableAttributes;\n";
+ os << formatv(" for (NamedAttribute attr : {0}) {{\n", attributes);
+ os << " StringRef name = attr.getName().getValue();\n";
+ os << " (void)name;\n";
+ os << " if (";
+ bool emittedName = false;
+ for (StringRef name : op.getInherentAttrNames()) {
+ if (emittedName)
+ os << " && ";
+ os << formatv("name != \"{0}\"", name);
+ emittedName = true;
+ }
+ os << ")\n";
+ os << " discardableAttributes.push_back(attr);\n";
+ os << " }\n";
+ os << formatv(" auto {1} = {0}::create(opBuilder, loc, {2}, {3}, "
+ "properties, discardableAttributes); (void){1};\n",
+ op.getQualCppClassName(), opVar, resultTypes, operands);
+ }
if (op.getNumResults() == 1) {
os << formatv(" valueMap[{0}] = {1}.getResult();\n\n", valueID, opVar);
}
diff --git a/mlir/tools/tblgen-to-irdl/OpDefinitionsGen.cpp b/mlir/tools/tblgen-to-irdl/OpDefinitionsGen.cpp
index 15418cac2fcc9..66a92c7bcf2e0 100644
--- a/mlir/tools/tblgen-to-irdl/OpDefinitionsGen.cpp
+++ b/mlir/tools/tblgen-to-irdl/OpDefinitionsGen.cpp
@@ -351,8 +351,9 @@ static Value createRegionConstraint(OpBuilder &builder,
if (predRec.getName() == "AnyRegion") {
ValueRange entryBlockArgs = {};
- auto op =
- irdl::RegionOp::create(builder, UnknownLoc::get(ctx), entryBlockArgs);
+ auto op = irdl::RegionOp::create(
+ builder, UnknownLoc::get(ctx), entryBlockArgs,
+ /*numberOfBlocks=*/nullptr, /*constrainedArguments=*/false);
return op.getResult();
}
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCOpsTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCOpsTest.cpp
index f277001ca1cfd..df1109d7f6dcf 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCOpsTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCOpsTest.cpp
@@ -19,6 +19,16 @@
using namespace mlir;
using namespace mlir::acc;
+template <typename OpTy>
+static OpTy createEmptyOp(OpBuilder &builder, Location loc) {
+ typename OpTy::Properties properties{};
+ OpTy::populateDefaultProperties(
+ OperationName(OpTy::getOperationName(), builder.getContext()),
+ properties);
+ return OpTy::create(builder, loc, TypeRange{}, ValueRange{}, properties,
+ /*discardableAttributes=*/{});
+}
+
//===----------------------------------------------------------------------===//
// Test Fixture
//===----------------------------------------------------------------------===//
@@ -45,7 +55,7 @@ class OpenACCOpsTest : public ::testing::Test {
template <typename Op>
static void testAsyncOnly(OpBuilder &b, MLIRContext &context, Location loc,
llvm::SmallVector<DeviceType> &dtypes) {
- OwningOpRef<Op> op = Op::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<Op> op = createEmptyOp<Op>(b, loc);
EXPECT_FALSE(op->hasAsyncOnly());
for (auto d : dtypes)
EXPECT_FALSE(op->hasAsyncOnly(d));
@@ -129,7 +139,7 @@ TEST_F(OpenACCOpsTest, asyncOnlyTestDataEntry) {
template <typename Op>
static void testAsyncValue(OpBuilder &b, MLIRContext &context, Location loc,
llvm::SmallVector<DeviceType> &dtypes) {
- OwningOpRef<Op> op = Op::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<Op> op = createEmptyOp<Op>(b, loc);
mlir::Value empty;
EXPECT_EQ(op->getAsyncValue(), empty);
@@ -200,7 +210,7 @@ static void
testNumGangsValues(OpBuilder &b, MLIRContext &context, Location loc,
llvm::SmallVector<DeviceType> &dtypes,
llvm::SmallVector<DeviceType> &dtypesWithoutNone) {
- OwningOpRef<Op> op = Op::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<Op> op = createEmptyOp<Op>(b, loc);
EXPECT_EQ(op->getNumGangsValues().begin(), op->getNumGangsValues().end());
OwningOpRef<arith::ConstantIndexOp> val1 =
@@ -267,7 +277,7 @@ TEST_F(OpenACCOpsTest, numGangsValuesTest) {
template <typename Op>
static void testVectorLength(OpBuilder &b, MLIRContext &context, Location loc,
llvm::SmallVector<DeviceType> &dtypes) {
- OwningOpRef<Op> op = Op::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<Op> op = createEmptyOp<Op>(b, loc);
mlir::Value empty;
EXPECT_EQ(op->getVectorLengthValue(), empty);
@@ -295,7 +305,7 @@ template <typename Op>
static void testWaitOnly(OpBuilder &b, MLIRContext &context, Location loc,
llvm::SmallVector<DeviceType> &dtypes,
llvm::SmallVector<DeviceType> &dtypesWithoutNone) {
- OwningOpRef<Op> op = Op::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<Op> op = createEmptyOp<Op>(b, loc);
EXPECT_FALSE(op->hasWaitOnly());
for (auto d : dtypes)
EXPECT_FALSE(op->hasWaitOnly(d));
@@ -335,7 +345,7 @@ template <typename Op>
static void testWaitValues(OpBuilder &b, MLIRContext &context, Location loc,
llvm::SmallVector<DeviceType> &dtypes,
llvm::SmallVector<DeviceType> &dtypesWithoutNone) {
- OwningOpRef<Op> op = Op::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<Op> op = createEmptyOp<Op>(b, loc);
EXPECT_EQ(op->getWaitValues().begin(), op->getWaitValues().end());
OwningOpRef<arith::ConstantIndexOp> val1 =
@@ -429,7 +439,7 @@ TEST_F(OpenACCOpsTest, waitValuesTest) {
}
TEST_F(OpenACCOpsTest, loopOpGangVectorWorkerTest) {
- OwningOpRef<LoopOp> op = LoopOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<LoopOp> op = createEmptyOp<LoopOp>(b, loc);
EXPECT_FALSE(op->hasGang());
EXPECT_FALSE(op->hasVector());
EXPECT_FALSE(op->hasWorker());
@@ -475,8 +485,7 @@ TEST_F(OpenACCOpsTest, loopOpGangVectorWorkerTest) {
}
TEST_F(OpenACCOpsTest, routineOpTest) {
- OwningOpRef<RoutineOp> op =
- RoutineOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<RoutineOp> op = createEmptyOp<RoutineOp>(b, loc);
EXPECT_FALSE(op->hasSeq());
EXPECT_FALSE(op->hasVector());
@@ -565,8 +574,7 @@ TEST_F(OpenACCOpsTest, routineOpTest) {
TEST_F(OpenACCOpsTest, routineOpGetBindNameValueOnlyBindStrOrOnlyBindId) {
// getBindNameValue(DeviceType) must not dereference when only one of bind(id)
// or bind(name) is set (the other has no device-type array).
- OwningOpRef<RoutineOp> op =
- RoutineOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<RoutineOp> op = createEmptyOp<RoutineOp>(b, loc);
auto dtypeNone = DeviceTypeAttr::get(&context, DeviceType::None);
auto dtypeNvidia = DeviceTypeAttr::get(&context, DeviceType::Nvidia);
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsLoopTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsLoopTest.cpp
index c7337f4967d4c..f13a09bbcaa8f 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsLoopTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsLoopTest.cpp
@@ -26,6 +26,16 @@
using namespace mlir;
using namespace mlir::acc;
+template <typename OpTy>
+static OpTy createEmptyOp(OpBuilder &builder, Location loc) {
+ typename OpTy::Properties properties{};
+ OpTy::populateDefaultProperties(
+ OperationName(OpTy::getOperationName(), builder.getContext()),
+ properties);
+ return OpTy::create(builder, loc, TypeRange{}, ValueRange{}, properties,
+ /*discardableAttributes=*/{});
+}
+
//===----------------------------------------------------------------------===//
// Test Fixture
//===----------------------------------------------------------------------===//
@@ -635,7 +645,7 @@ TEST_F(OpenACCUtilsLoopTest,
// only to own the region, then build entry -> then/else -> exit with
// acc.yield.
OwningOpRef<acc::ParallelOp> parallelOp =
- acc::ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ createEmptyOp<acc::ParallelOp>(b, loc);
Region ®ion = parallelOp->getRegion();
Block *entry = b.createBlock(®ion, region.begin());
Block *thenBlock = b.createBlock(®ion, region.end());
@@ -700,7 +710,7 @@ TEST_F(OpenACCUtilsLoopTest,
auto [module, funcOp] = createModuleWithFunc();
OwningOpRef<acc::ParallelOp> parallelOp =
- acc::ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ createEmptyOp<acc::ParallelOp>(b, loc);
Region ®ion = parallelOp->getRegion();
// Block order as in all.mlir: ^bb0 entry, ^bb1 header, ^bb2 exit, ^bb3 body
Block *entry = b.createBlock(®ion, region.begin());
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTest.cpp
index 4f519135e5111..bd0be8abd0da9 100644
--- a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTest.cpp
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTest.cpp
@@ -23,6 +23,16 @@
using namespace mlir;
using namespace mlir::acc;
+template <typename OpTy>
+static OpTy createEmptyOp(OpBuilder &builder, Location loc) {
+ typename OpTy::Properties properties{};
+ OpTy::populateDefaultProperties(
+ OperationName(OpTy::getOperationName(), builder.getContext()),
+ properties);
+ return OpTy::create(builder, loc, TypeRange{}, ValueRange{}, properties,
+ /*discardableAttributes=*/{});
+}
+
//===----------------------------------------------------------------------===//
// Test Fixture
//===----------------------------------------------------------------------===//
@@ -46,8 +56,7 @@ class OpenACCUtilsTest : public ::testing::Test {
TEST_F(OpenACCUtilsTest, getEnclosingComputeOpParallel) {
// Create a parallel op with a region
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
Region ¶llelRegion = parallelOp->getRegion();
parallelRegion.emplaceBlock();
@@ -58,8 +67,7 @@ TEST_F(OpenACCUtilsTest, getEnclosingComputeOpParallel) {
TEST_F(OpenACCUtilsTest, getEnclosingComputeOpKernels) {
// Create a kernels op with a region
- OwningOpRef<KernelsOp> kernelsOp =
- KernelsOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<KernelsOp> kernelsOp = createEmptyOp<KernelsOp>(b, loc);
Region &kernelsRegion = kernelsOp->getRegion();
kernelsRegion.emplaceBlock();
@@ -70,8 +78,7 @@ TEST_F(OpenACCUtilsTest, getEnclosingComputeOpKernels) {
TEST_F(OpenACCUtilsTest, getEnclosingComputeOpSerial) {
// Create a serial op with a region
- OwningOpRef<SerialOp> serialOp =
- SerialOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<SerialOp> serialOp = createEmptyOp<SerialOp>(b, loc);
Region &serialRegion = serialOp->getRegion();
serialRegion.emplaceBlock();
@@ -82,8 +89,7 @@ TEST_F(OpenACCUtilsTest, getEnclosingComputeOpSerial) {
TEST_F(OpenACCUtilsTest, getEnclosingComputeOpNested) {
// Create nested ops: parallel containing a loop op
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
Region ¶llelRegion = parallelOp->getRegion();
Block *parallelBlock = ¶llelRegion.emplaceBlock();
@@ -91,8 +97,7 @@ TEST_F(OpenACCUtilsTest, getEnclosingComputeOpNested) {
b.setInsertionPointToStart(parallelBlock);
// Create a loop op inside the parallel region
- OwningOpRef<LoopOp> loopOp =
- LoopOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<LoopOp> loopOp = createEmptyOp<LoopOp>(b, loc);
Region &loopRegion = loopOp->getRegion();
loopRegion.emplaceBlock();
@@ -125,8 +130,7 @@ TEST_F(OpenACCUtilsTest, isOnlyUsedByPrivateClausesTrue) {
cast<TypedValue<PointerLikeType>>(allocOp->getResult());
// Create a parallel op with a region
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
Region ¶llelRegion = parallelOp->getRegion();
Block *parallelBlock = ¶llelRegion.emplaceBlock();
@@ -150,8 +154,7 @@ TEST_F(OpenACCUtilsTest, isOnlyUsedByPrivateClausesFalse) {
cast<TypedValue<PointerLikeType>>(allocOp->getResult());
// Create a parallel op with a region
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
Region ¶llelRegion = parallelOp->getRegion();
Block *parallelBlock = ¶llelRegion.emplaceBlock();
@@ -180,8 +183,7 @@ TEST_F(OpenACCUtilsTest, isOnlyUsedByPrivateClausesMultiple) {
cast<TypedValue<PointerLikeType>>(allocOp->getResult());
// Create a parallel op with a region
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
Region ¶llelRegion = parallelOp->getRegion();
Block *parallelBlock = ¶llelRegion.emplaceBlock();
@@ -211,8 +213,7 @@ TEST_F(OpenACCUtilsTest, isOnlyUsedByReductionClausesTrue) {
cast<TypedValue<PointerLikeType>>(allocOp->getResult());
// Create a parallel op with a region
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
Region ¶llelRegion = parallelOp->getRegion();
Block *parallelBlock = ¶llelRegion.emplaceBlock();
@@ -236,8 +237,7 @@ TEST_F(OpenACCUtilsTest, isOnlyUsedByReductionClausesFalse) {
cast<TypedValue<PointerLikeType>>(allocOp->getResult());
// Create a parallel op with a region
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
Region ¶llelRegion = parallelOp->getRegion();
Block *parallelBlock = ¶llelRegion.emplaceBlock();
@@ -266,8 +266,7 @@ TEST_F(OpenACCUtilsTest, isOnlyUsedByReductionClausesMultiple) {
cast<TypedValue<PointerLikeType>>(allocOp->getResult());
// Create a parallel op with a region
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
Region ¶llelRegion = parallelOp->getRegion();
Block *parallelBlock = ¶llelRegion.emplaceBlock();
@@ -291,8 +290,7 @@ TEST_F(OpenACCUtilsTest, isOnlyUsedByReductionClausesMultiple) {
TEST_F(OpenACCUtilsTest, getDefaultAttrOnParallel) {
// Create a parallel op with a default attribute
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
parallelOp->setDefaultAttr(ClauseDefaultValue::None);
// Test that we can retrieve the default attribute
@@ -304,8 +302,7 @@ TEST_F(OpenACCUtilsTest, getDefaultAttrOnParallel) {
TEST_F(OpenACCUtilsTest, getDefaultAttrOnKernels) {
// Create a kernels op with a default attribute
- OwningOpRef<KernelsOp> kernelsOp =
- KernelsOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<KernelsOp> kernelsOp = createEmptyOp<KernelsOp>(b, loc);
kernelsOp->setDefaultAttr(ClauseDefaultValue::Present);
// Test that we can retrieve the default attribute
@@ -317,8 +314,7 @@ TEST_F(OpenACCUtilsTest, getDefaultAttrOnKernels) {
TEST_F(OpenACCUtilsTest, getDefaultAttrOnSerial) {
// Create a serial op with a default attribute
- OwningOpRef<SerialOp> serialOp =
- SerialOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<SerialOp> serialOp = createEmptyOp<SerialOp>(b, loc);
serialOp->setDefaultAttr(ClauseDefaultValue::None);
// Test that we can retrieve the default attribute
@@ -330,8 +326,7 @@ TEST_F(OpenACCUtilsTest, getDefaultAttrOnSerial) {
TEST_F(OpenACCUtilsTest, getDefaultAttrOnData) {
// Create a data op with a default attribute
- OwningOpRef<DataOp> dataOp =
- DataOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<DataOp> dataOp = createEmptyOp<DataOp>(b, loc);
dataOp->setDefaultAttr(ClauseDefaultValue::Present);
// Test that we can retrieve the default attribute
@@ -342,8 +337,7 @@ TEST_F(OpenACCUtilsTest, getDefaultAttrOnData) {
TEST_F(OpenACCUtilsTest, getDefaultAttrNone) {
// Create a parallel op without setting a default attribute
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
// Do not set default attribute
// Test that we get std::nullopt when there's no default attribute
@@ -354,8 +348,7 @@ TEST_F(OpenACCUtilsTest, getDefaultAttrNone) {
TEST_F(OpenACCUtilsTest, getDefaultAttrNearest) {
// Create a data op with a default attribute
- OwningOpRef<DataOp> dataOp =
- DataOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<DataOp> dataOp = createEmptyOp<DataOp>(b, loc);
dataOp->setDefaultAttr(ClauseDefaultValue::Present);
Region &dataRegion = dataOp->getRegion();
@@ -365,8 +358,7 @@ TEST_F(OpenACCUtilsTest, getDefaultAttrNearest) {
b.setInsertionPointToStart(dataBlock);
// Create a parallel op inside the data region with NO default attribute
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
// Do not set default attribute on parallel op
Region ¶llelRegion = parallelOp->getRegion();
@@ -375,8 +367,7 @@ TEST_F(OpenACCUtilsTest, getDefaultAttrNearest) {
b.setInsertionPointToStart(parallelBlock);
// Create a loop op inside the parallel region
- OwningOpRef<LoopOp> loopOp =
- LoopOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<LoopOp> loopOp = createEmptyOp<LoopOp>(b, loc);
// Test that from the loop op, we find the nearest default attribute (from
// data op)
@@ -975,8 +966,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesFromComputeConstruct) {
/*name=*/"test_var");
// Create a parallel op
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
// Set the data clause operands
parallelOp->getDataClauseOperandsMutable().append(copyinOp->getAccVar());
@@ -1023,8 +1013,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesFromEnclosingDataOp) {
/*name=*/"test_var");
// Create a data op
- OwningOpRef<DataOp> dataOp =
- DataOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<DataOp> dataOp = createEmptyOp<DataOp>(b, loc);
// Set the data clause operands
dataOp->getDataClauseOperandsMutable().append(copyinOp->getAccVar());
@@ -1035,8 +1024,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesFromEnclosingDataOp) {
b.setInsertionPointToStart(dataBlock);
// Create a parallel op inside the data region (no data clauses on parallel)
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
// Create dominance info
DominanceInfo domInfo(funcOp.get());
@@ -1088,8 +1076,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesFromComputeAndEnclosingData) {
/*name=*/"var2");
// Create a data op
- OwningOpRef<DataOp> dataOp =
- DataOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<DataOp> dataOp = createEmptyOp<DataOp>(b, loc);
// Set the data clause operands for data op
dataOp->getDataClauseOperandsMutable().append(copyinOp1->getAccVar());
@@ -1100,8 +1087,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesFromComputeAndEnclosingData) {
b.setInsertionPointToStart(dataBlock);
// Create a parallel op inside the data region
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
// Set the data clause operands for parallel op
parallelOp->getDataClauseOperandsMutable().append(copyinOp2->getAccVar());
@@ -1155,8 +1141,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesWithDeclareDirectives) {
ValueRange{copyinOp->getAccVar()});
// Create a parallel op
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
// Create a declare_exit op that post-dominates the parallel
OwningOpRef<DeclareExitOp> declareExitOp = DeclareExitOp::create(
@@ -1223,8 +1208,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesMultipleDataConstructs) {
/*name=*/"var3");
// Create outer data op
- OwningOpRef<DataOp> outerDataOp =
- DataOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<DataOp> outerDataOp = createEmptyOp<DataOp>(b, loc);
// Set the data clause operands for outer data op
outerDataOp->getDataClauseOperandsMutable().append(copyinOp1->getAccVar());
@@ -1235,8 +1219,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesMultipleDataConstructs) {
b.setInsertionPointToStart(outerDataBlock);
// Create inner data op
- OwningOpRef<DataOp> innerDataOp =
- DataOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<DataOp> innerDataOp = createEmptyOp<DataOp>(b, loc);
// Set the data clause operands for inner data op
innerDataOp->getDataClauseOperandsMutable().append(copyinOp2->getAccVar());
@@ -1247,8 +1230,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesMultipleDataConstructs) {
b.setInsertionPointToStart(innerDataBlock);
// Create a parallel op
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
// Set the data clause operands for parallel op
parallelOp->getDataClauseOperandsMutable().append(copyinOp3->getAccVar());
@@ -1297,8 +1279,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesKernelsOp) {
/*name=*/"test_var");
// Create a kernels op
- OwningOpRef<KernelsOp> kernelsOp =
- KernelsOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<KernelsOp> kernelsOp = createEmptyOp<KernelsOp>(b, loc);
// Set the data clause operands
kernelsOp->getDataClauseOperandsMutable().append(copyinOp->getAccVar());
@@ -1345,8 +1326,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesSerialOp) {
/*name=*/"test_var");
// Create a serial op
- OwningOpRef<SerialOp> serialOp =
- SerialOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<SerialOp> serialOp = createEmptyOp<SerialOp>(b, loc);
// Set the data clause operands
serialOp->getDataClauseOperandsMutable().append(copyinOp->getAccVar());
@@ -1381,8 +1361,7 @@ TEST_F(OpenACCUtilsTest, getDominatingDataClausesEmpty) {
b.setInsertionPointToStart(funcBlock);
// Create a parallel op with no data clauses
- OwningOpRef<ParallelOp> parallelOp =
- ParallelOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<ParallelOp> parallelOp = createEmptyOp<ParallelOp>(b, loc);
// Create dominance info
DominanceInfo domInfo(funcOp.get());
@@ -1607,8 +1586,7 @@ TEST_F(OpenACCUtilsTest, isValidValueUseFromDataEntryOp) {
Value dataClauseResult = copyinOp->getAccVar();
// Create a serial region
- OwningOpRef<SerialOp> serialOp =
- SerialOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<SerialOp> serialOp = createEmptyOp<SerialOp>(b, loc);
Region &serialRegion = serialOp->getRegion();
// Value from data entry op should be valid
@@ -1640,8 +1618,7 @@ TEST_F(OpenACCUtilsTest, isValidValueUseDeviceData) {
Value deviceVal = allocOp->getResult();
// Create a serial region
- OwningOpRef<SerialOp> serialOp =
- SerialOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<SerialOp> serialOp = createEmptyOp<SerialOp>(b, loc);
Region &serialRegion = serialOp->getRegion();
// Device data should be valid
@@ -1671,8 +1648,7 @@ TEST_F(OpenACCUtilsTest, isValidValueUseOnlyUsedByPrivate) {
cast<TypedValue<PointerLikeType>>(allocOp->getResult());
// Create a serial region with a private clause using the variable
- OwningOpRef<SerialOp> serialOp =
- SerialOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<SerialOp> serialOp = createEmptyOp<SerialOp>(b, loc);
Region &serialRegion = serialOp->getRegion();
Block *serialBlock = b.createBlock(&serialRegion);
b.setInsertionPointToStart(serialBlock);
@@ -1706,8 +1682,7 @@ TEST_F(OpenACCUtilsTest, isValidValueUseRegularValue) {
Value regularVal = allocOp->getResult();
// Create a serial region with a non-private use of the value
- OwningOpRef<SerialOp> serialOp =
- SerialOp::create(b, loc, TypeRange{}, ValueRange{});
+ OwningOpRef<SerialOp> serialOp = createEmptyOp<SerialOp>(b, loc);
Region &serialRegion = serialOp->getRegion();
Block *serialBlock = b.createBlock(&serialRegion);
b.setInsertionPointToStart(serialBlock);
diff --git a/mlir/unittests/TableGen/OpBuildGen.cpp b/mlir/unittests/TableGen/OpBuildGen.cpp
index e069f4a07de27..9d69f0fd32068 100644
--- a/mlir/unittests/TableGen/OpBuildGen.cpp
+++ b/mlir/unittests/TableGen/OpBuildGen.cpp
@@ -290,13 +290,17 @@ TEST_F(OpBuildGenTest, BuildMethodsVariadicProperties) {
verifyOp(std::move(op), {f32Ty}, {*cstI32}, {*cstI32}, noAttrs);
// Test build method with result types, supplied attributes.
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_PUSH
op = test::TableGenBuildOp6::create(builder, loc, TypeRange{f32Ty},
ValueRange{*cstI32, *cstI32}, attrs);
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_POP
verifyOp(std::move(op), {f32Ty}, {*cstI32}, {*cstI32}, attrs);
// Test build method with no result types and supplied attributes.
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_PUSH
op = test::TableGenBuildOp6::create(builder, loc,
ValueRange{*cstI32, *cstI32}, attrs);
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_POP
verifyOp(std::move(op), {f32Ty}, {*cstI32}, {*cstI32}, attrs);
// Test replacing an inherent attribute backed by a native property.
@@ -335,8 +339,10 @@ TEST_F(OpBuildGenTest, BuildMethodsInherentDiscardableAttrs) {
// Check that the old-style builder partitions the attributes and populates
// properties before Operation::create.
OperationState state(loc, test::TableGenBuildOp7::getOperationName());
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_PUSH
test::TableGenBuildOp7::build(builder, state, TypeRange{}, ValueRange{},
attrs);
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_POP
ASSERT_TRUE(state.getRawProperties());
EXPECT_EQ(state.attributes.getAttrs().size(), 1u);
EXPECT_EQ(state.attributes.getAttrs()[0], attrs[1]);
@@ -347,9 +353,11 @@ TEST_F(OpBuildGenTest, BuildMethodsInherentDiscardableAttrs) {
auto op7FromState = cast<test::TableGenBuildOp7>(builder.create(state));
verifyOp(op7FromState, {}, {}, attrs);
- // Check that the legacy create forwarder remains compatible.
+ // Check that the deprecated create forwarder remains compatible.
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_PUSH
auto op7b = test::TableGenBuildOp7::create(builder, loc, TypeRange{},
ValueRange{}, attrs);
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_POP
// Note: this goes before verifyOp() because verifyOp() calls erase(), causing
// use-after-free.
ASSERT_EQ(op7b.getProperties().getAttr0(), attrs[0].getValue());
@@ -367,8 +375,10 @@ TEST_F(OpBuildGenTest, BuildMethodsLegacyMixedProperties) {
builder.getNamedAttr("unknown", builder.getStringAttr("discardable"))};
OperationState state(loc, test::TableGenBuildOp8::getOperationName());
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_PUSH
test::TableGenBuildOp8::build(builder, state, ValueRange{*cstI32, *cstF32},
mixedAttrs);
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_POP
ASSERT_TRUE(state.getRawProperties());
ASSERT_EQ(state.attributes.getAttrs().size(), 1u);
@@ -396,8 +406,10 @@ TEST_F(OpBuildGenTest, BuildMethodsLegacySameOperandAndResultType) {
SmallVector<NamedAttribute> mixedAttrs{
builder.getNamedAttr("attr0", builder.getBoolAttr(true)),
builder.getNamedAttr("unknown", builder.getUnitAttr())};
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_PUSH
auto op = test::TableGenBuildOp9::create(builder, loc, ValueRange{*cstI32},
mixedAttrs);
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_POP
EXPECT_EQ(op.getResult().getType(), i32Ty);
EXPECT_TRUE(op.getAttr0());
EXPECT_EQ(op->getDiscardableAttrDictionary().size(), 1u);
@@ -409,8 +421,10 @@ TEST_F(OpBuildGenTest, BuildMethodsLegacyFirstAttrDerivedResultType) {
SmallVector<NamedAttribute> mixedAttrs{
builder.getNamedAttr("type", TypeAttr::get(f32Ty)),
builder.getNamedAttr("unknown", builder.getUnitAttr())};
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_PUSH
auto op = test::TableGenBuildOp10::create(builder, loc, ValueRange{*cstI32},
mixedAttrs);
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_POP
EXPECT_EQ(op.getResult().getType(), f32Ty);
EXPECT_EQ(op.getType(), f32Ty);
EXPECT_EQ(op->getDiscardableAttrDictionary().size(), 1u);
@@ -430,10 +444,12 @@ TEST_F(OpBuildGenTest, BuildMethodsInvalidLegacyPropertyConversion) {
SmallVector<NamedAttribute> badAttrs{
builder.getNamedAttr("attr0", builder.getStringAttr("not-a-bool"))};
OperationState state(loc, test::TableGenBuildOp7::getOperationName());
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_PUSH
EXPECT_DEATH_IF_SUPPORTED(
test::TableGenBuildOp7::build(builder, state, TypeRange{}, ValueRange{},
badAttrs),
"Invalid attribute.*attr0");
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_POP
}
} // namespace mlir
More information about the llvm-branch-commits
mailing list