[flang-commits] [flang] fc696f7 - [mlir] Migrate aggregate builders to explicit properties (#219195)
via flang-commits
flang-commits at lists.llvm.org
Thu Sep 10 13:53:36 PDT 2026
Author: Mehdi Amini
Date: 2026-09-10T22:53:30+02:00
New Revision: fc696f791e06bdc0c0aa10afa65e2243db528409
URL: https://github.com/llvm/llvm-project/commit/fc696f791e06bdc0c0aa10afa65e2243db528409
DIFF: https://github.com/llvm/llvm-project/commit/fc696f791e06bdc0c0aa10afa65e2243db528409.diff
LOG: [mlir] Migrate aggregate builders to explicit properties (#219195)
Pass typed property structs and discardable attributes separately at ODS
aggregate builder call sites.
This is part of a general migration to use the "new" properties-based
APIs #155475
Assisted-by: Codex
Added:
flang/test/Fir/call-properties.fir
mlir/test/Conversion/FuncToLLVM/call-attributes.mlir
Modified:
flang/lib/Lower/OpenACC.cpp
flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp
flang/lib/Optimizer/Builder/IntrinsicCall.cpp
flang/lib/Optimizer/Builder/PPCIntrinsicCall.cpp
flang/lib/Optimizer/CodeGen/CodeGen.cpp
flang/lib/Optimizer/CodeGen/TargetRewrite.cpp
mlir/docs/DefiningDialects/Operations.md
mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h
mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td
mlir/include/mlir/Dialect/LLVMIR/LLVMOpBase.td
mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td
mlir/include/mlir/Dialect/Tensor/IR/TensorOps.td
mlir/include/mlir/TableGen/Operator.h
mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
mlir/lib/Conversion/ArithToAMDGPU/ArithToAMDGPU.cpp
mlir/lib/Conversion/ArithToEmitC/ArithToEmitC.cpp
mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp
mlir/lib/Conversion/MathToFuncs/MathToFuncs.cpp
mlir/lib/Conversion/MathToLLVM/MathToLLVM.cpp
mlir/lib/Conversion/MathToLibm/MathToLibm.cpp
mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
mlir/lib/Conversion/RaiseWasm/RaiseWasmMLIR.cpp
mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp
mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp
mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp
mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp
mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp
mlir/lib/Conversion/VectorToSCF/VectorToSCF.cpp
mlir/lib/Conversion/VectorToXeGPU/VectorToXeGPU.cpp
mlir/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp
mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
mlir/lib/Dialect/Arith/IR/ArithOps.cpp
mlir/lib/Dialect/Arith/Transforms/EmulateUnsupportedFloats.cpp
mlir/lib/Dialect/Arith/Transforms/ExpandOps.cpp
mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp
mlir/lib/Dialect/Arith/Utils/Utils.cpp
mlir/lib/Dialect/GPU/Transforms/DecomposeMemRefs.cpp
mlir/lib/Dialect/Linalg/Transforms/Vectorization.cpp
mlir/lib/Dialect/Math/Transforms/ExtendToSupportedTypes.cpp
mlir/lib/Dialect/Math/Transforms/PolynomialApproximation.cpp
mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp
mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp
mlir/lib/Dialect/OpenACC/Transforms/ACCCGToGPU.cpp
mlir/lib/Dialect/Shape/IR/Shape.cpp
mlir/lib/Dialect/SparseTensor/Transforms/SparseVectorization.cpp
mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp
mlir/lib/Dialect/Tensor/IR/TensorOps.cpp
mlir/lib/Dialect/Tosa/Transforms/TosaNarrowTypes.cpp
mlir/lib/Dialect/Vector/Transforms/LowerVectorContract.cpp
mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
mlir/lib/TableGen/Operator.cpp
mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp
mlir/lib/Tools/PDLL/CodeGen/MLIRGen.cpp
mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
mlir/test/Dialect/SPIRV/IR/target-env.mlir
mlir/test/Dialect/Tosa/tosa-narrow-i64-to-i32.mlir
mlir/test/Target/SPIRV/memory-ops.mlir
mlir/test/lib/Dialect/SPIRV/TestAvailability.cpp
mlir/test/lib/Dialect/Test/TestPatterns.cpp
mlir/test/mlir-tblgen/op-decl-and-defs.td
mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp
mlir/tools/mlir-tblgen/SPIRVUtilsGen.cpp
mlir/tools/tblgen-to-irdl/OpDefinitionsGen.cpp
mlir/unittests/Dialect/OpenACC/OpenACCOpsTest.cpp
mlir/unittests/Dialect/OpenACC/OpenACCUtilsLoopTest.cpp
mlir/unittests/Dialect/OpenACC/OpenACCUtilsTest.cpp
mlir/unittests/TableGen/OpBuildGen.cpp
Removed:
################################################################################
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 8d1499b260fa8..d185260126abe 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..6bbe65731964f 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,9 @@ 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,
+ /*alias_scopes=*/{},
+ /*noalias_scopes=*/{}, /*tbaa=*/{})
.getResult();
break;
}
@@ -2262,7 +2264,9 @@ 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,
+ /*alias_scopes=*/{},
+ /*noalias_scopes=*/{}, /*tbaa=*/{})
.getResult();
break;
}
@@ -2270,15 +2274,18 @@ 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,
+ /*alias_scopes=*/{},
+ /*noalias_scopes=*/{}, /*tbaa=*/{})
.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,
+ /*alias_scopes=*/{}, /*noalias_scopes=*/{}, /*tbaa=*/{})
.getResult();
break;
}
@@ -2290,32 +2297,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,
+ /*alias_scopes=*/{}, /*noalias_scopes=*/{}, /*tbaa=*/{})
.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,
+ /*alias_scopes=*/{}, /*noalias_scopes=*/{}, /*tbaa=*/{})
.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,
+ /*alias_scopes=*/{}, /*noalias_scopes=*/{}, /*tbaa=*/{})
.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,
+ /*alias_scopes=*/{}, /*noalias_scopes=*/{}, /*tbaa=*/{})
.getResult();
break;
}
@@ -2377,8 +2388,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 +2398,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
diff ers 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 19dc4142d645d..d782665951708 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