[flang-commits] [flang] [mlir] [mlir] Migrate aggregate builders to explicit properties (PR #219195)
Mehdi Amini via flang-commits
flang-commits at lists.llvm.org
Thu Sep 10 11:05:47 PDT 2026
https://github.com/joker-eph updated https://github.com/llvm/llvm-project/pull/219195
>From 4864170cbc4799dc30be63ed5b8579a2c1e3f2dd 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. Preserve direct FIR callees and cover call
property and discardable attribute preservation in FIR and Func lowering.
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 | 111 +++++++++++++----
flang/lib/Optimizer/CodeGen/TargetRewrite.cpp | 3 +-
flang/test/Fir/call-properties.fir | 21 ++++
mlir/docs/DefiningDialects/Operations.md | 18 +--
.../ArithCommon/AttrToLLVMConverter.h | 40 ++++--
.../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 | 55 ++++----
.../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 | 63 ++++++----
mlir/lib/Conversion/MathToLibm/MathToLibm.cpp | 13 +-
mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp | 13 +-
.../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 | 46 +++++--
mlir/lib/Dialect/Arith/IR/ArithOps.cpp | 10 +-
.../Transforms/EmulateUnsupportedFloats.cpp | 3 +-
.../Dialect/Arith/Transforms/ExpandOps.cpp | 3 +-
.../Transforms/UnsignedWhenEquivalent.cpp | 6 +-
mlir/lib/Dialect/Arith/Utils/Utils.cpp | 11 +-
.../GPU/Transforms/DecomposeMemRefs.cpp | 2 +-
.../Linalg/Transforms/Vectorization.cpp | 8 +-
.../Transforms/ExtendToSupportedTypes.cpp | 4 +-
.../Transforms/PolynomialApproximation.cpp | 9 +-
mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp | 14 ++-
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 +-
.../Transforms/XeGPUWgToSgDistribute.cpp | 9 +-
mlir/lib/TableGen/Operator.cpp | 29 +++++
.../SPIRV/Deserialization/DeserializeOps.cpp | 13 +-
mlir/lib/Tools/PDLL/CodeGen/MLIRGen.cpp | 3 +-
.../FuncToLLVM/call-attributes.mlir | 17 +++
.../SPIRVToLLVM/constant-op-to-llvm.mlir | 7 ++
mlir/test/Dialect/SPIRV/IR/target-env.mlir | 3 +-
.../Dialect/Tosa/tosa-narrow-i64-to-i32.mlir | 4 +-
mlir/test/Target/SPIRV/memory-ops.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 | 73 +++++------
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 | 117 +++++++-----------
mlir/unittests/TableGen/OpBuildGen.cpp | 20 ++-
70 files changed, 932 insertions(+), 440 deletions(-)
create mode 100644 flang/test/Fir/call-properties.fir
create mode 100644 mlir/test/Conversion/FuncToLLVM/call-attributes.mlir
diff --git a/flang/lib/Lower/OpenACC.cpp b/flang/lib/Lower/OpenACC.cpp
index d883478e49375..42e5fa98326ce 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;
@@ -1304,7 +1305,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 a7267d834cf6a..db41d0f5151bd 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 a139af9a29e57..ddd083ef029e8 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 67be3201b82db..1fafee2e90714 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,16 @@ 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.getDiscardableAttrs(),
+ adaptor.getOperands().size());
+ builderAttrs.properties.callee =
+ mlir::cast_if_present<mlir::FlatSymbolRefAttr>(call.getCalleeAttr());
+ 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 +1536,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 +1561,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 +1647,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();
}
@@ -3624,10 +3679,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();
}
@@ -3994,9 +4050,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/flang/test/Fir/call-properties.fir b/flang/test/Fir/call-properties.fir
new file mode 100644
index 0000000000000..3c3971a0bcf91
--- /dev/null
+++ b/flang/test/Fir/call-properties.fir
@@ -0,0 +1,21 @@
+// RUN: fir-opt --fir-to-llvm-ir="target=x86_64-unknown-linux-gnu" %s | FileCheck %s
+
+// Preserve the direct callee alongside converted fast-math properties and
+// discardable attributes. An indirect call must keep its function operand.
+func.func private @callee(f32) -> f32
+
+func.func @direct(%arg: f32) -> f32 {
+ %result = fir.call @callee(%arg) fastmath<contract> {test.tag = "direct"} : (f32) -> f32
+ return %result : f32
+}
+
+// CHECK-LABEL: llvm.func @direct(
+// CHECK: llvm.call @callee({{.*}}) {{.*}}fastmathFlags = #llvm.fastmath<contract>{{.*}}test.tag = "direct"
+
+func.func @indirect(%callee: (f32) -> f32, %arg: f32) -> f32 {
+ %result = fir.call %callee(%arg) fastmath<contract> {test.tag = "indirect"} : (f32) -> f32
+ return %result : f32
+}
+
+// CHECK-LABEL: llvm.func @indirect(
+// CHECK: llvm.call %{{.*}}({{.*}}) {{.*}}fastmathFlags = #llvm.fastmath<contract>{{.*}}test.tag = "indirect"
diff --git a/mlir/docs/DefiningDialects/Operations.md b/mlir/docs/DefiningDialects/Operations.md
index faf84811f1fb4..5615af678d339 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 cdea5307aa72d..b76290c8176be 100644
--- a/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h
+++ b/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h
@@ -46,27 +46,47 @@ convertArithRoundingModeAttrToLLVM(arith::RoundingModeAttr roundingModeAttr);
LLVM::FPExceptionBehaviorAttr
getLLVMDefaultFPExceptionBehavior(MLIRContext &context);
-// Attribute converter that populates a NamedAttrList by removing the fastmath
-// attribute from the source operation attributes, and replacing it with an
-// equivalent LLVM fastmath attribute.
+// Convert the source fastmath attribute to LLVM properties, keeping the source
+// discardable attributes separate.
template <typename SourceOp, typename TargetOp>
class AttrConvertFastMathToLLVM {
public:
AttrConvertFastMathToLLVM(SourceOp srcOp)
- : convertedAttr(srcOp->getDiscardableAttrDictionary()) {
- srcOp->getName().populateInherentAttrs(srcOp, convertedAttr);
- convertedAttr.erase(SourceOp::getFastMathAttrName());
+ : context(srcOp.getOperation()->getContext()),
+ convertedAttr(srcOp->getDiscardableAttrDictionary()) {
auto arithFMFAttr = srcOp.getFastMathFlagsAttr();
if (arithFMFAttr) {
- convertedAttr.set(TargetOp::getFastmathAttrName(),
- convertArithFastMathAttrToLLVM(arithFMFAttr));
+ StringRef targetAttrName = TargetOp::getFastmathAttrName();
+ Builder builder(context);
+ propertiesAttr = builder.getDictionaryAttr(builder.getNamedAttr(
+ targetAttrName, convertArithFastMathAttrToLLVM(arithFMFAttr)));
}
}
- ArrayRef<NamedAttribute> getAttrs() const { return convertedAttr.getAttrs(); }
- Attribute getPropAttr() const { return {}; }
+ ArrayRef<NamedAttribute> getDiscardableAttrs() const {
+ return convertedAttr.getAttrs();
+ }
+ ArrayRef<NamedAttribute> getAttrs() const { return getDiscardableAttrs(); }
+ 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 eb46921319e88..9ab067af36f79 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
@@ -523,8 +523,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 c71f9d0a3f54f..34fb1c6b3153c 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 341a70554a9f3..23acce19f89e5 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);
@@ -2254,7 +2254,8 @@ struct TransposeLoadOpLowering
if (numElements != 16)
return emitNumElementsError(16, "gfx1250+");
intrinsic =
- ROCDL::DsLoadTr4_B64::create(rewriter, loc, rocdlResultType, srcPtr)
+ ROCDL::DsLoadTr4_B64::create(rewriter, loc, rocdlResultType, srcPtr,
+ ArrayAttr{}, ArrayAttr{}, ArrayAttr{})
.getResult();
break;
}
@@ -2262,7 +2263,8 @@ struct TransposeLoadOpLowering
if (numElements != 16)
return emitNumElementsError(16, "gfx1250+");
intrinsic =
- ROCDL::DsLoadTr6_B96::create(rewriter, loc, rocdlResultType, srcPtr)
+ ROCDL::DsLoadTr6_B96::create(rewriter, loc, rocdlResultType, srcPtr,
+ ArrayAttr{}, ArrayAttr{}, ArrayAttr{})
.getResult();
break;
}
@@ -2270,15 +2272,17 @@ struct TransposeLoadOpLowering
if (numElements != 8)
return emitNumElementsError(8, "gfx1250+");
intrinsic =
- ROCDL::DsLoadTr8_B64::create(rewriter, loc, rocdlResultType, srcPtr)
+ ROCDL::DsLoadTr8_B64::create(rewriter, loc, rocdlResultType, srcPtr,
+ ArrayAttr{}, ArrayAttr{}, ArrayAttr{})
.getResult();
break;
}
case 16: {
if (numElements != 8)
return emitNumElementsError(8, "gfx1250+");
- intrinsic = ROCDL::DsLoadTr16_B128::create(rewriter, loc,
- rocdlResultType, srcPtr)
+ intrinsic = ROCDL::DsLoadTr16_B128::create(
+ rewriter, loc, rocdlResultType, srcPtr, ArrayAttr{},
+ ArrayAttr{}, ArrayAttr{})
.getResult();
break;
}
@@ -2290,32 +2294,36 @@ struct TransposeLoadOpLowering
case 4: {
if (numElements != 16)
return emitNumElementsError(16, "gfx950");
- intrinsic = ROCDL::ds_read_tr4_b64::create(rewriter, loc,
- rocdlResultType, srcPtr)
+ intrinsic = ROCDL::ds_read_tr4_b64::create(
+ rewriter, loc, rocdlResultType, srcPtr, ArrayAttr{},
+ ArrayAttr{}, ArrayAttr{})
.getResult();
break;
}
case 6: {
if (numElements != 16)
return emitNumElementsError(16, "gfx950");
- intrinsic = ROCDL::ds_read_tr6_b96::create(rewriter, loc,
- rocdlResultType, srcPtr)
+ intrinsic = ROCDL::ds_read_tr6_b96::create(
+ rewriter, loc, rocdlResultType, srcPtr, ArrayAttr{},
+ ArrayAttr{}, ArrayAttr{})
.getResult();
break;
}
case 8: {
if (numElements != 8)
return emitNumElementsError(8, "gfx950");
- intrinsic = ROCDL::ds_read_tr8_b64::create(rewriter, loc,
- rocdlResultType, srcPtr)
+ intrinsic = ROCDL::ds_read_tr8_b64::create(
+ rewriter, loc, rocdlResultType, srcPtr, ArrayAttr{},
+ ArrayAttr{}, ArrayAttr{})
.getResult();
break;
}
case 16: {
if (numElements != 4)
return emitNumElementsError(4, "gfx950");
- intrinsic = ROCDL::ds_read_tr16_b64::create(rewriter, loc,
- rocdlResultType, srcPtr)
+ intrinsic = ROCDL::ds_read_tr16_b64::create(
+ rewriter, loc, rocdlResultType, srcPtr, ArrayAttr{},
+ ArrayAttr{}, ArrayAttr{})
.getResult();
break;
}
@@ -2377,8 +2385,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, rocdlResultType, srcPtr, ArrayAttr{}, ArrayAttr{},
+ ArrayAttr{});
rewriter.replaceOpWithNewOp<LLVM::BitcastOp>(op, llvmResultType, rocdlOp);
break;
}
@@ -2386,22 +2395,24 @@ 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, rocdlResultType, srcPtr, ArrayAttr{}, ArrayAttr{},
+ ArrayAttr{});
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, rocdlResultType, srcPtr, ArrayAttr{}, ArrayAttr{},
+ ArrayAttr{});
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, llvmResultType, srcPtr, ArrayAttr{}, ArrayAttr{}, ArrayAttr{});
break;
}
default:
diff --git a/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp b/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp
index 6d99475283b8e..d0714739589f5 100644
--- a/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp
+++ b/mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp
@@ -111,7 +111,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");
}
@@ -150,8 +151,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);
@@ -202,7 +204,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");
@@ -471,7 +475,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);
@@ -589,7 +595,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 268b8de5cba9e..7567e61228013 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..29e1678bdd38a 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.getDiscardableAttrs());
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.getDiscardableAttrs());
rewriter.replaceOpWithNewOp<LLVM::FSubOp>(
- op, llvmOperandType, ValueRange{exp, one}, subAttrs.getAttrs());
+ op, TypeRange{llvmOperandType}, ValueRange{exp, one},
+ subAttrs.getProperties(), subAttrs.getDiscardableAttrs());
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.getDiscardableAttrs());
+ return LLVM::FSubOp::create(
+ rewriter, loc, TypeRange{llvm1DVectorTy}, ValueRange{exp, one},
+ subAttrs.getProperties(), subAttrs.getDiscardableAttrs());
},
rewriter);
}
@@ -297,11 +302,13 @@ struct Log1pOpLowering
: LLVM::ConstantOp::create(rewriter, loc, llvmOperandType,
floatOne);
- auto add = LLVM::FAddOp::create(rewriter, loc, llvmOperandType,
+ auto add = LLVM::FAddOp::create(rewriter, loc, TypeRange{llvmOperandType},
ValueRange{one, adaptor.getOperand()},
- addAttrs.getAttrs());
+ addAttrs.getProperties(),
+ addAttrs.getDiscardableAttrs());
rewriter.replaceOpWithNewOp<LLVM::LogOp>(
- op, llvmOperandType, ValueRange{add}, logAttrs.getAttrs());
+ op, TypeRange{llvmOperandType}, ValueRange{add},
+ logAttrs.getProperties(), logAttrs.getDiscardableAttrs());
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.getDiscardableAttrs());
+ return LLVM::LogOp::create(rewriter, loc, TypeRange{llvm1DVectorTy},
+ ValueRange{add}, logAttrs.getProperties(),
+ logAttrs.getDiscardableAttrs());
},
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.getDiscardableAttrs());
rewriter.replaceOpWithNewOp<LLVM::FDivOp>(
- op, llvmOperandType, ValueRange{one, sqrt}, divAttrs.getAttrs());
+ op, TypeRange{llvmOperandType}, ValueRange{one, sqrt},
+ divAttrs.getProperties(), divAttrs.getDiscardableAttrs());
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.getDiscardableAttrs());
+ return LLVM::FDivOp::create(
+ rewriter, loc, TypeRange{llvm1DVectorTy}, ValueRange{one, sqrt},
+ divAttrs.getProperties(), divAttrs.getDiscardableAttrs());
},
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 58b7948892db4..f96d67aa64071 100644
--- a/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
+++ b/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
@@ -144,16 +144,9 @@ 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);
- SmallVector<NamedAttribute> discardableAttrs;
- for (mlir::NamedAttribute attr : fastAttrConverter.getAttrs()) {
- if (attr.getName() == LLVM::CallOp::getFastmathAttrName()) {
- callOp.setFastmathFlagsAttr(
- cast<LLVM::FastmathFlagsAttr>(attr.getValue()));
- continue;
- }
- discardableAttrs.push_back(attr);
- }
- callOp->setDiscardableAttrs(discardableAttrs);
+ callOp.setFastmathFlagsAttr(
+ fastAttrConverter.getProperties().getFastmathFlags());
+ callOp->setDiscardableAttrs(fastAttrConverter.getDiscardableAttrs());
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 fe73f72812f24..7b8303d87db68 100644
--- a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
+++ b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
@@ -864,7 +864,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 971239a6caeac..03eb6dcccb380 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 cebf40bbe2c38..6b98c242aeac3 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 7c227d4f49106..b8eed6315e389 100644
--- a/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp
+++ b/mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp
@@ -619,7 +619,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();
}
@@ -662,7 +663,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->setDiscardableAttr(kPassLabel, rewriter.getUnitAttr());
diff --git a/mlir/lib/Conversion/VectorToXeGPU/VectorToXeGPU.cpp b/mlir/lib/Conversion/VectorToXeGPU/VectorToXeGPU.cpp
index 706f24128720b..3a56d21240109 100644
--- a/mlir/lib/Conversion/VectorToXeGPU/VectorToXeGPU.cpp
+++ b/mlir/lib/Conversion/VectorToXeGPU/VectorToXeGPU.cpp
@@ -966,9 +966,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 662f83ce2702e..cab0e53380ca3 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 b59f7e97553c1..ce47da8d9ee97 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,19 +126,25 @@ LogicalResult RawBufferAtomicByCasPattern<AtomicOp, ArithOp>::matchAndRewrite(
ConversionPatternRewriter &rewriter) const {
Location loc = atomicOp.getLoc();
- NamedAttrList origAttrs(atomicOp->getDiscardableAttrDictionary());
- atomicOp->getName().walkInherentAttrs(
- atomicOp,
- [&](StringRef name, Attribute &attr) { origAttrs.append(name, attr); });
+ NamedAttrList origProperties;
+ atomicOp->getName().walkInherentAttrs(atomicOp,
+ [&](StringRef name, Attribute &attr) {
+ origProperties.append(name, attr);
+ });
+ 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());
@@ -138,11 +160,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 58ea71f6ebf7e..95ec15e515eeb 100644
--- a/mlir/lib/Dialect/Arith/IR/ArithOps.cpp
+++ b/mlir/lib/Dialect/Arith/IR/ArithOps.cpp
@@ -1396,12 +1396,10 @@ struct NarrowExtremum final : OpRewritePattern<TruncOp> {
return failure();
}
- OperationState state(truncOp.getLoc(), ExtremumOp::getOperationName(),
- ValueRange{lhs, rhs}, TypeRange{narrowType},
- extremumOp->getDiscardableAttrDictionary().getValue());
- state.propertiesAttr = extremumOp->getPropertiesAsAttribute();
- Operation *newExtremum = rewriter.create(state);
- rewriter.replaceOp(truncOp, newExtremum->getResults());
+ 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 23659d8010f45..9b336cbda62f5 100644
--- a/mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp
+++ b/mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp
@@ -100,7 +100,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 ab23515c02ee9..a6fabe42d1e42 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 a1bd74d5cb1a1..aac591bd8125d 100644
--- a/mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp
+++ b/mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp
@@ -92,8 +92,12 @@ struct ConvertOpToUnsigned final : OpRewritePattern<Signed> {
staticallyNonNegative(this->solver, static_cast<Operation *>(op))))
return failure();
+ typename Unsigned::Properties properties{};
+ Unsigned::populateDefaultProperties(
+ OperationName(Unsigned::getOperationName(), rw.getContext()),
+ properties);
rw.replaceOpWithNewOp<Unsigned>(
- op, op->getResultTypes(), op->getOperands(),
+ 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 e19f59f0e4478..21887d46759c8 100644
--- a/mlir/lib/Dialect/Linalg/Transforms/Vectorization.cpp
+++ b/mlir/lib/Dialect/Linalg/Transforms/Vectorization.cpp
@@ -3721,10 +3721,10 @@ struct Conv1DGenerator
else
dstType = dstElementType;
- return rewriter
- .create(loc, castOp->getName().getIdentifier(), val, dstType,
- castOp->getAttrs())
- ->getResult(0);
+ OperationState state(loc, castOp->getName().getIdentifier(), val, dstType,
+ castOp->getDiscardableAttrDictionary().getValue());
+ state.propertiesAttr = castOp->getPropertiesAsAttribute();
+ return rewriter.create(state)->getResult(0);
}
// Create a contraction: lhs{n, w, c} * rhs{c, f} -> res{n, w, f}
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/OpenACC/IR/OpenACC.cpp b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
index 332226846f4dd..0857513f7f7a0 100644
--- a/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
+++ b/mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
@@ -225,7 +225,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 6a46810a56a1c..b9d634e059a01 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
@@ -2776,11 +2776,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);
@@ -3948,7 +3948,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);
@@ -3993,7 +3994,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()),
@@ -4009,7 +4011,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 c4f85bea14b49..43bb5cabb0f4f 100644
--- a/mlir/lib/Dialect/Tensor/IR/TensorOps.cpp
+++ b/mlir/lib/Dialect/Tensor/IR/TensorOps.cpp
@@ -2056,9 +2056,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 570d96e0995fa..ba8ac48053e9a 100644
--- a/mlir/lib/Dialect/Tosa/Transforms/TosaNarrowTypes.cpp
+++ b/mlir/lib/Dialect/Tosa/Transforms/TosaNarrowTypes.cpp
@@ -574,7 +574,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..2a152ef0634d7 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,
+ /*fastmath=*/{});
return arith::ExtSIOp::create(rewriter, loc, promotedType, v);
}
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
index accc401b573bd..8169d2a282f38 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
@@ -322,12 +322,9 @@ struct WgToSgDpasOp : public OpConversionPattern<xegpu::DpasOp> {
for (auto aVec : adaptor.getLhs()) {
for (auto bVec : adaptor.getRhs()) {
- llvm::SmallVector<Value> operands({aVec, bVec});
Value tmpC;
- if (op.getAcc()) {
+ if (op.getAcc())
tmpC = adaptor.getAcc()[i++];
- operands.push_back(tmpC);
- }
ArrayRef<int64_t> aVecShape =
cast<VectorType>(aVec.getType()).getShape();
@@ -339,7 +336,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, aVec, bVec, 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 a3cb63f4f3fa4..f7f92ab69238f 100644
--- a/mlir/lib/Tools/PDLL/CodeGen/MLIRGen.cpp
+++ b/mlir/lib/Tools/PDLL/CodeGen/MLIRGen.cpp
@@ -442,7 +442,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/Conversion/FuncToLLVM/call-attributes.mlir b/mlir/test/Conversion/FuncToLLVM/call-attributes.mlir
new file mode 100644
index 0000000000000..518b0f356aad2
--- /dev/null
+++ b/mlir/test/Conversion/FuncToLLVM/call-attributes.mlir
@@ -0,0 +1,17 @@
+// RUN: mlir-opt -convert-func-to-llvm %s | FileCheck %s
+
+func.func private @callee(f32) -> f32
+
+// CHECK-LABEL: llvm.func @direct_call
+func.func @direct_call(%arg: f32) -> f32 {
+ // CHECK: llvm.call fastcc @callee(%arg0) {convergent, fastmathFlags = #llvm.fastmath<fast>, test.marker = "keep"}
+ %0 = func.call @callee(%arg) {CConv = #llvm.cconv<fastcc>, fastmathFlags = #llvm.fastmath<fast>, convergent, test.marker = "keep"} : (f32) -> f32
+ return %0 : f32
+}
+
+// CHECK-LABEL: llvm.func @indirect_call
+func.func @indirect_call(%fn: (f32) -> f32, %arg: f32) -> f32 {
+ // CHECK: llvm.call fastcc %arg0(%arg1) {convergent, fastmathFlags = #llvm.fastmath<nnan>, test.marker = "keep"}
+ %0 = func.call_indirect %fn(%arg) {CConv = #llvm.cconv<fastcc>, fastmathFlags = #llvm.fastmath<nnan>, convergent, test.marker = "keep"} : (f32) -> f32
+ return %0 : f32
+}
diff --git a/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
index 2d74022b34406..ae4f1d9cbe27b 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
+++ b/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
@@ -59,3 +59,10 @@ spirv.func @float_constant_vector() "None" {
%0 = spirv.Constant dense<[2.000000e+00, 3.000000e+00]> : vector<2xf32>
spirv.Return
}
+
+// CHECK-LABEL: @constant_with_discardable_attr
+spirv.func @constant_with_discardable_attr() "None" {
+ // CHECK: llvm.mlir.constant(42 : i32) {test.marker = "keep"} : i32
+ %0 = "spirv.Constant"() <{value = 42 : i32}> {test.marker = "keep"} : () -> i32
+ spirv.Return
+}
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/Target/SPIRV/memory-ops.mlir b/mlir/test/Target/SPIRV/memory-ops.mlir
index fa786f18aac18..91c5ccdf78464 100644
--- a/mlir/test/Target/SPIRV/memory-ops.mlir
+++ b/mlir/test/Target/SPIRV/memory-ops.mlir
@@ -18,8 +18,8 @@ spirv.module Logical GLSL450 requires #spirv.vce<v1.0, [Shader, Linkage], []> {
// CHECK-LABEL: spirv.func @load_store_memory_operands
spirv.func @load_store_memory_operands(%arg0 : !spirv.ptr<f32, Input>, %arg1 : !spirv.ptr<f32, Output>) "None" {
- // CHECK: spirv.Load "Input" %{{.+}} ["Volatile|Aligned", 4] : f32
- %1 = spirv.Load "Input" %arg0 ["Volatile|Aligned", 4]: f32
+ // CHECK: spirv.Load "Input" %{{.+}} ["Volatile|Aligned", 4] {relaxed_precision} : f32
+ %1 = spirv.Load "Input" %arg0 ["Volatile|Aligned", 4] {relaxed_precision} : f32
// CHECK: spirv.Store "Output" %{{.+}}, %{{.+}} ["Volatile|Aligned", 4] : f32
spirv.Store "Output" %arg1, %1 ["Volatile|Aligned", 4]: f32
spirv.Return
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 06517910e270e..5fc7c1e801741 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 8becd36353cdf..2b195ce7e31cf 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";
-
/// The logic to calculate the actual value range for a declared operand/result
/// of an op with variadic operands/results. Note that this logic is not for
/// general use; it assumes all variadic operands/results must have the same
@@ -382,16 +375,7 @@ class OpOrAdaptorHelper {
/// Returns whether the operation will have a non-empty `Properties` struct.
bool hasNonEmptyPropertiesStruct() const {
- if (!op.getProperties().empty())
- return true;
- 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 op.hasNonEmptyProperties();
}
std::optional<NamedProperty> &getOperandSegmentsSize() {
@@ -626,7 +610,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
@@ -2418,7 +2403,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;
@@ -2449,6 +2434,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);
@@ -2588,21 +2579,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();
auto *method = opClass.addStaticMethod<Method::Private>(
"void", "buildPropertiesAndDiscardableAttributes",
@@ -2685,7 +2662,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
@@ -2759,7 +2740,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();
@@ -2863,7 +2848,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();
@@ -3029,7 +3018,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 6a86465311439..0e25cf9ca4ebf 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 d50a62ebaa538..e83a26a184cb5 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)
@@ -996,8 +987,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());
@@ -1044,8 +1034,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());
@@ -1056,8 +1045,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());
@@ -1109,8 +1097,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());
@@ -1121,8 +1108,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());
@@ -1176,8 +1162,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(
@@ -1244,8 +1229,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());
@@ -1256,8 +1240,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());
@@ -1268,8 +1251,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());
@@ -1318,8 +1300,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());
@@ -1366,8 +1347,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());
@@ -1402,8 +1382,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());
@@ -1628,8 +1607,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
@@ -1661,8 +1639,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
@@ -1692,8 +1669,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);
@@ -1727,8 +1703,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);
@@ -1758,7 +1733,7 @@ TEST_F(OpenACCUtilsTest, isValidValueUseRoutineArgument) {
Block *entryBlock = funcOp.addEntryBlock();
b.setInsertionPointToStart(entryBlock);
- auto serialOp = SerialOp::create(b, loc, TypeRange{}, ValueRange{});
+ auto serialOp = createEmptyOp<SerialOp>(b, loc);
Block *serialBlock = b.createBlock(&serialOp.getRegion());
b.setInsertionPointToStart(serialBlock);
func::CallOp::create(b, loc, "use_arg", TypeRange{},
diff --git a/mlir/unittests/TableGen/OpBuildGen.cpp b/mlir/unittests/TableGen/OpBuildGen.cpp
index b053051982e17..e0e1ac99884ba 100644
--- a/mlir/unittests/TableGen/OpBuildGen.cpp
+++ b/mlir/unittests/TableGen/OpBuildGen.cpp
@@ -291,13 +291,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.
@@ -336,8 +340,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]);
@@ -348,9 +354,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());
@@ -368,8 +376,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);
@@ -395,9 +405,11 @@ TEST_F(OpBuildGenTest, BuildMethodsLegacyMixedProperties) {
TEST_F(OpBuildGenTest, BuildMethodsLegacyDefaultsWithoutAttributes) {
OperationState state(loc, test::TableGenBuildOp8::getOperationName());
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_PUSH
test::TableGenBuildOp8::build(builder, state, TypeRange{i32Ty},
ValueRange{*cstI32, *cstF32},
ArrayRef<NamedAttribute>{});
+ LLVM_SUPPRESS_DEPRECATED_DECLARATIONS_POP
ASSERT_TRUE(state.getRawProperties());
const auto &properties =
@@ -409,8 +421,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);
@@ -422,8 +436,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);
@@ -443,10 +459,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 flang-commits
mailing list