[Mlir-commits] [mlir] [mlir][LLVM][GPU] Migrate to explicit split inherent/discardable attribute APIs access (PR #218921)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Wed Aug 26 06:14:30 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-amdgpu
@llvm/pr-subscribers-mlir-ods
Author: Mehdi Amini (joker-eph)
<details>
<summary>Changes</summary>
Use discardable attribute APIs and typed operation accessors throughout the LLVM and GPU dialect families, their conversions, translations, and tests.
Assisted-by: Codex
---
Patch is 86.69 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/218921.diff
56 Files Affected:
- (modified) mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h (+15-36)
- (modified) mlir/include/mlir/Conversion/LLVMCommon/Pattern.h (+3-2)
- (modified) mlir/include/mlir/Conversion/LLVMCommon/VectorPattern.h (+5-2)
- (modified) mlir/include/mlir/Dialect/GPU/IR/GPUOps.td (+4-2)
- (modified) mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td (+12-5)
- (modified) mlir/include/mlir/Dialect/LLVMIR/LLVMInterfaces.td (+11)
- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+2-1)
- (modified) mlir/include/mlir/IR/SymbolInterfaces.td (+4-3)
- (modified) mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp (+3-3)
- (modified) mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp (+4-4)
- (modified) mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp (+2-2)
- (modified) mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp (+2-2)
- (modified) mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp (+8-6)
- (modified) mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h (+1-1)
- (modified) mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp (+2-2)
- (modified) mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp (+6-5)
- (modified) mlir/lib/Conversion/LLVMCommon/LowerFunctionDiscardablesToLLVM.cpp (+2-1)
- (modified) mlir/lib/Conversion/LLVMCommon/Pattern.cpp (+1-1)
- (modified) mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp (+2-1)
- (modified) mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp (+2-2)
- (modified) mlir/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp (+2-2)
- (modified) mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp (+23-9)
- (modified) mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp (+4-1)
- (modified) mlir/lib/Dialect/AMDGPU/Transforms/MaskedloadToLoad.cpp (+2-2)
- (modified) mlir/lib/Dialect/GPU/IR/GPUDialect.cpp (+26-18)
- (modified) mlir/lib/Dialect/GPU/IR/InferIntRangeInterfaceImpls.cpp (+3-1)
- (modified) mlir/lib/Dialect/GPU/Transforms/EliminateBarriers.cpp (+1-1)
- (modified) mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp (+5-3)
- (modified) mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp (+4-3)
- (modified) mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp (+2-2)
- (modified) mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp (+19-14)
- (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+16-11)
- (modified) mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp (+1-1)
- (modified) mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp (+3-2)
- (modified) mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp (+2-3)
- (modified) mlir/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp (+1-1)
- (modified) mlir/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp (+11-9)
- (modified) mlir/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp (+5-5)
- (modified) mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp (+6-4)
- (modified) mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp (+8-5)
- (modified) mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp (+14-16)
- (modified) mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp (+3-2)
- (modified) mlir/lib/Dialect/XeGPU/Utils/XeGPUUtils.cpp (+15-12)
- (modified) mlir/lib/Target/LLVM/XeVM/Target.cpp (+2-1)
- (modified) mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMIRToLLVMTranslation.cpp (+1-1)
- (modified) mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp (+1-1)
- (modified) mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp (+4-4)
- (modified) mlir/lib/Target/LLVMIR/ModuleImport.cpp (+15-12)
- (modified) mlir/lib/Target/LLVMIR/ModuleTranslation.cpp (+7-5)
- (modified) mlir/lib/Target/LLVMIR/Transforms/TargetToDataLayout.cpp (+5-4)
- (modified) mlir/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp (+3-2)
- (modified) mlir/python/mlir/dialects/gpu/__init__.py (+3-7)
- (added) mlir/test/Conversion/ArithToLLVM/attribute-storage.mlir (+32)
- (modified) mlir/test/lib/Dialect/LLVM/TestPatterns.cpp (+3-2)
- (modified) mlir/test/lib/Dialect/XeGPU/TestXeGPUTransforms.cpp (+3-2)
- (modified) mlir/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp (+3-2)
``````````diff
diff --git a/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h b/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h
index feb74c86e349f..2a4e489715d02 100644
--- a/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h
+++ b/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h
@@ -18,6 +18,7 @@
namespace mlir {
namespace arith {
+
/// Maps arithmetic fastmath enum values to LLVM enum values.
LLVM::FastmathFlags
convertArithFastMathFlagsToLLVM(arith::FastMathFlags arithFMF);
@@ -51,17 +52,11 @@ getLLVMDefaultFPExceptionBehavior(MLIRContext &context);
template <typename SourceOp, typename TargetOp>
class AttrConvertFastMathToLLVM {
public:
- AttrConvertFastMathToLLVM(SourceOp srcOp) {
- // Copy the source attributes.
- convertedAttr = NamedAttrList{srcOp->getAttrs()};
- // Get the name of the arith fastmath attribute.
- StringRef arithFMFAttrName = SourceOp::getFastMathAttrName();
- // Remove the source fastmath attribute.
- auto arithFMFAttr = dyn_cast_if_present<arith::FastMathFlagsAttr>(
- convertedAttr.erase(arithFMFAttrName));
+ AttrConvertFastMathToLLVM(SourceOp srcOp)
+ : convertedAttr(srcOp->getDiscardableAttrDictionary()) {
+ auto arithFMFAttr = srcOp.getFastMathFlagsAttr();
if (arithFMFAttr) {
- StringRef targetAttrName = TargetOp::getFastmathAttrName();
- convertedAttr.set(targetAttrName,
+ convertedAttr.set(TargetOp::getFastmathAttrName(),
convertArithFastMathAttrToLLVM(arithFMFAttr));
}
}
@@ -78,17 +73,11 @@ class AttrConvertFastMathToLLVM {
template <typename SourceOp, typename TargetOp>
class AttrConvertOverflowToLLVM {
public:
- AttrConvertOverflowToLLVM(SourceOp srcOp) {
+ AttrConvertOverflowToLLVM(SourceOp srcOp)
+ : convertedAttr(srcOp->getDiscardableAttrDictionary()) {
using IntegerOverflowFlagsAttr = LLVM::IntegerOverflowFlagsAttr;
- // Copy the source attributes.
- convertedAttr = NamedAttrList{srcOp->getAttrs()};
- // Get the name of the arith overflow attribute.
- StringRef arithAttrName = SourceOp::getIntegerOverflowAttrName();
- // Remove the source overflow attribute from the set that will be present
- // in the target.
- if (auto arithAttr = dyn_cast_if_present<arith::IntegerOverflowFlagsAttr>(
- convertedAttr.erase(arithAttrName))) {
+ if (auto arithAttr = srcOp.getOverflowAttr()) {
auto llvmFlag = convertArithOverflowFlagsToLLVM(arithAttr.getValue());
// Create a dictionary attribute holding the overflow flags property.
// (In the LLVM dialect, the overflow flags are a property, not an
@@ -117,9 +106,9 @@ class AttrConvertOverflowToLLVM {
template <typename SourceOp, typename TargetOp>
class AttrConvertNonNegToLLVM {
public:
- AttrConvertNonNegToLLVM(SourceOp srcOp) {
- convertedAttr = NamedAttrList{srcOp->getAttrs()};
- if (!convertedAttr.erase("nonNeg"))
+ AttrConvertNonNegToLLVM(SourceOp srcOp)
+ : convertedAttr(srcOp->getDiscardableAttrDictionary()) {
+ if (!srcOp.getNonNeg())
return;
MLIRContext *ctx = srcOp.getOperation()->getContext();
Builder b(ctx);
@@ -142,26 +131,16 @@ class AttrConverterConstrainedFPToLLVM {
"LLVM::FPExceptionBehaviorOpInterface");
public:
- AttrConverterConstrainedFPToLLVM(SourceOp srcOp) {
- // Copy the source attributes.
- convertedAttr = NamedAttrList{srcOp->getAttrs()};
-
+ AttrConverterConstrainedFPToLLVM(SourceOp srcOp)
+ : convertedAttr(srcOp->getDiscardableAttrDictionary()) {
if constexpr (TargetOp::template hasTrait<
LLVM::RoundingModeOpInterface::Trait>()) {
- // Get the name of the rounding mode attribute.
- StringRef arithAttrName = srcOp.getRoundingModeAttrName();
- // Remove the source attribute.
- auto arithAttr =
- cast<arith::RoundingModeAttr>(convertedAttr.erase(arithAttrName));
- // Set the target attribute.
+ auto arithAttr = srcOp.getRoundingModeAttr();
convertedAttr.set(TargetOp::getRoundingModeAttrName(),
convertArithRoundingModeAttrToLLVM(arithAttr));
}
// Constrained intrinsics (llvm.intr.experimental.constrained.*) do not
- // support fastmath flags. Remove the arith fastmath attribute if present.
- if constexpr (SourceOp::template hasTrait<
- arith::ArithFastMathInterface::Trait>())
- convertedAttr.erase(srcOp.getFastMathAttrName());
+ // support fastmath flags, so do not copy them from the source operation.
convertedAttr.set(TargetOp::getFPExceptionBehaviorAttrName(),
getLLVMDefaultFPExceptionBehavior(*srcOp->getContext()));
}
diff --git a/mlir/include/mlir/Conversion/LLVMCommon/Pattern.h b/mlir/include/mlir/Conversion/LLVMCommon/Pattern.h
index 2f468458addd3..b50aeb1e274b3 100644
--- a/mlir/include/mlir/Conversion/LLVMCommon/Pattern.h
+++ b/mlir/include/mlir/Conversion/LLVMCommon/Pattern.h
@@ -338,8 +338,9 @@ class OneToOneConvertToLLVMPattern : public ConvertOpToLLVMPattern<SourceOp> {
matchAndRewrite(SourceOp op, typename SourceOp::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
return LLVM::detail::oneToOneRewrite(
- op, TargetOp::getOperationName(), adaptor.getOperands(), op->getAttrs(),
- /*propertiesAttr=*/Attribute{}, *this->getTypeConverter(), rewriter);
+ op, TargetOp::getOperationName(), adaptor.getOperands(),
+ op->getDiscardableAttrDictionary().getValue(),
+ op->getPropertiesAsAttribute(), *this->getTypeConverter(), rewriter);
}
};
diff --git a/mlir/include/mlir/Conversion/LLVMCommon/VectorPattern.h b/mlir/include/mlir/Conversion/LLVMCommon/VectorPattern.h
index 65988a2466318..fdc787d84383e 100644
--- a/mlir/include/mlir/Conversion/LLVMCommon/VectorPattern.h
+++ b/mlir/include/mlir/Conversion/LLVMCommon/VectorPattern.h
@@ -70,13 +70,16 @@ LogicalResult vectorOneToOneRewrite(Operation *op, StringRef targetOp,
template <typename SourceOp, typename TargetOp>
class AttrConvertPassThrough {
public:
- AttrConvertPassThrough(SourceOp srcOp) : srcAttrs(srcOp->getAttrs()) {}
+ AttrConvertPassThrough(SourceOp srcOp)
+ : srcAttrs(srcOp->getDiscardableAttrDictionary().getValue()),
+ propertiesAttr(srcOp->getPropertiesAsAttribute()) {}
ArrayRef<NamedAttribute> getAttrs() const { return srcAttrs; }
- Attribute getPropAttr() const { return {}; }
+ Attribute getPropAttr() const { return propertiesAttr; }
private:
ArrayRef<NamedAttribute> srcAttrs;
+ Attribute propertiesAttr;
};
/// Basic lowering implementation to rewrite Ops with just one result to the
diff --git a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
index 1066dcb3b7308..34cdf019d5966 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -419,7 +419,9 @@ def GPU_GPUFuncOp : GPU_Op<"func", [
attribution.
}];
- let arguments = (ins TypeAttrOf<FunctionType>:$function_type,
+ let arguments = (ins SymbolNameAttr:$sym_name,
+ OptionalAttr<StrAttr>:$sym_visibility,
+ TypeAttrOf<FunctionType>:$function_type,
OptionalAttr<DictArrayAttr>:$arg_attrs,
OptionalAttr<DictArrayAttr>:$res_attrs,
OptionalAttr<DictArrayAttr>:$workgroup_attrib_attrs,
@@ -447,7 +449,7 @@ def GPU_GPUFuncOp : GPU_Op<"func", [
bool isKernel() {
if (getKernel())
return true;
- return (*this)->getAttrOfType<UnitAttr>(
+ return (*this)->getDiscardableAttrOfType<UnitAttr>(
GPUDialect::getKernelFuncAttrName()) != nullptr;
}
diff --git a/mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td b/mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td
index 2834ae160595d..b631ec169f943 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td
@@ -157,11 +157,18 @@ def BasicPtxBuilderOpInterface : OpInterface<"BasicPtxBuilderInterface"> {
for (auto val : op->getOperands())
asmValues.push_back({val, mlir::NVVM::PTXRegisterMod::Read});
- // Step 3. Add attributes
- for (auto attr : op->getAttrs()) {
- if (auto intAttr = dyn_cast<mlir::IntegerAttr>(attr.getValue())) {
- ::mlir::Value val = makeConstantI32(rewriter, intAttr.getInt());
- asmValues.push_back({val, mlir::NVVM::PTXRegisterMod::Read});
+ // Step 3. Add inherent attributes.
+ auto properties = llvm::dyn_cast_or_null<mlir::DictionaryAttr>(
+ op->getPropertiesAsAttribute());
+ if (properties) {
+ for (auto attr : properties) {
+ if (auto intAttr =
+ dyn_cast<mlir::IntegerAttr>(attr.getValue())) {
+ ::mlir::Value val =
+ makeConstantI32(rewriter, intAttr.getInt());
+ asmValues.push_back(
+ {val, mlir::NVVM::PTXRegisterMod::Read});
+ }
}
}
return false; // No manual mapping needed
diff --git a/mlir/include/mlir/Dialect/LLVMIR/LLVMInterfaces.td b/mlir/include/mlir/Dialect/LLVMIR/LLVMInterfaces.td
index c59b38c416956..66a1fe0ce251a 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/LLVMInterfaces.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/LLVMInterfaces.td
@@ -35,6 +35,17 @@ def FastmathFlagsInterface : OpInterface<"FastmathFlagsInterface"> {
return op.getFastmathFlagsAttr();
}]
>,
+ InterfaceMethod<
+ /*desc=*/ "Sets the FastmathFlagsAttr attribute for the operation",
+ /*returnType=*/ "void",
+ /*methodName=*/ "setFastmathAttr",
+ /*args=*/ (ins "::mlir::LLVM::FastmathFlagsAttr":$attr),
+ /*methodBody=*/ [{}],
+ /*defaultImpl=*/ [{
+ auto op = cast<ConcreteOp>(this->getOperation());
+ op.setFastmathFlagsAttr(attr);
+ }]
+ >,
StaticInterfaceMethod<
/*desc=*/ [{Returns the name of the FastmathFlagsAttr attribute
for the operation}],
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 6ac489f9ad2bc..795df0433dcf0 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -279,7 +279,8 @@ class NVVM_PureSpecialRangeableRegisterOp<string mnemonic, list<Trait> traits =
void $cppClass::inferResultRanges(
ArrayRef<::mlir::ConstantIntRanges> argRanges,
SetIntRangeFn setResultRanges) {
- nvvmInferResultRanges(getOperation(), getResult(), argRanges, setResultRanges);
+ nvvmInferResultRanges(getRange(), getResult(), argRanges,
+ setResultRanges);
}
// Verify the range attribute satisfies LLVM ConstantRange constructor requirements.
diff --git a/mlir/include/mlir/IR/SymbolInterfaces.td b/mlir/include/mlir/IR/SymbolInterfaces.td
index ebe0c26637ad3..5aad305e8c35a 100644
--- a/mlir/include/mlir/IR/SymbolInterfaces.td
+++ b/mlir/include/mlir/IR/SymbolInterfaces.td
@@ -42,8 +42,7 @@ def Symbol : OpInterface<"SymbolOpInterface"> {
InterfaceMethod<"Sets the name of this symbol.",
"void", "setName", (ins "::mlir::StringAttr":$name), [{}],
/*defaultImplementation=*/[{
- this->getOperation()->setAttr(
- mlir::SymbolTable::getSymbolAttrName(), name);
+ mlir::SymbolTable::setSymbolName(this->getOperation(), name);
}]
>,
InterfaceMethod<"Gets the visibility of this symbol.",
@@ -197,7 +196,9 @@ def Symbol : OpInterface<"SymbolOpInterface"> {
// Add additional classof checks to properly handle "optional" symbols.
let extraClassOf = [{
- return $_op->hasAttr(::mlir::SymbolTable::getSymbolAttrName());
+ return static_cast<bool>(
+ $_op->getInherentAttr(::mlir::SymbolTable::getSymbolAttrName())
+ .value_or(::mlir::Attribute{}));
}];
}
diff --git a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
index f47d095fd86e6..898dc8ff580c3 100644
--- a/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
+++ b/mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
@@ -3218,9 +3218,9 @@ struct AMDGPUDPPLowering : public ConvertOpToLLVMPattern<DPPOp> {
// Check for row_mask, bank_mask, bound_ctrl if they exist and create
// constants
- auto rowMask = DppOp->getAttrOfType<IntegerAttr>("row_mask").getInt();
- auto bankMask = DppOp->getAttrOfType<IntegerAttr>("bank_mask").getInt();
- bool boundCtrl = DppOp->getAttrOfType<BoolAttr>("bound_ctrl").getValue();
+ auto rowMask = DppOp.getRowMask();
+ auto bankMask = DppOp.getBankMask();
+ bool boundCtrl = DppOp.getBoundCtrl();
// create a ROCDL_DPPMovOp instruction with the appropriate attributes
auto dppMovOp =
diff --git a/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp b/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
index 4e8ef252a10c9..d766c47a817f5 100644
--- a/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
+++ b/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
@@ -374,10 +374,10 @@ struct SelectOpOneToNLowering : public ConvertOpToLLVMPattern<arith::SelectOp> {
LogicalResult
ConstantOpLowering::matchAndRewrite(arith::ConstantOp op, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const {
- return LLVM::detail::oneToOneRewrite(op, LLVM::ConstantOp::getOperationName(),
- adaptor.getOperands(), op->getAttrs(),
- /*propAttr=*/Attribute{},
- *getTypeConverter(), rewriter);
+ return LLVM::detail::oneToOneRewrite(
+ op, LLVM::ConstantOp::getOperationName(), adaptor.getOperands(),
+ op->getDiscardableAttrDictionary().getValue(),
+ op->getPropertiesAsAttribute(), *getTypeConverter(), rewriter);
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp b/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp
index ceda3888360ce..74acbe325c656 100644
--- a/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp
+++ b/mlir/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp
@@ -96,8 +96,8 @@ struct ConstantOpLowering : public ConvertOpToLLVMPattern<complex::ConstantOp> {
ConversionPatternRewriter &rewriter) const override {
return LLVM::detail::oneToOneRewrite(
op, LLVM::ConstantOp::getOperationName(), adaptor.getOperands(),
- op->getAttrs(), /*propAttr=*/Attribute{}, *getTypeConverter(),
- rewriter);
+ op->getDiscardableAttrDictionary().getValue(),
+ op->getPropertiesAsAttribute(), *getTypeConverter(), rewriter);
}
};
diff --git a/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp b/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
index fef78a46d69fc..bf82d30d527ad 100644
--- a/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
+++ b/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
@@ -148,12 +148,12 @@ struct BranchOpLowering : public ConvertOpToLLVMPattern<cf::BranchOp> {
TypeRange(ValueRange(flattenedAdaptor)));
if (failed(convertedBlock))
return failure();
- DictionaryAttr attrs = op->getAttrDictionary();
+ DictionaryAttr attrs = op->getDiscardableAttrDictionary();
Operation *newOp = rewriter.replaceOpWithNewOp<LLVM::BrOp>(
op, flattenedAdaptor, *convertedBlock);
// TODO: We should not just forward all attributes like that. But there are
// existing Flang tests that depend on this behavior.
- newOp->setAttrs(attrs);
+ newOp->setDiscardableAttrs(attrs);
return success();
}
};
diff --git a/mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp b/mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp
index eeb90a7ff8150..3ce209213f2a0 100644
--- a/mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp
+++ b/mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp
@@ -630,7 +630,6 @@ static Value scalarizeVectorOpHelper(Operation *op, ValueRange operands,
Location loc = op->getLoc();
Value result = LLVM::PoisonOp::create(rewriter, loc, vectorType);
Type indexType = converter.convertType(rewriter.getIndexType());
- StringAttr name = op->getName().getIdentifier();
Type elementType = vectorType.getElementType();
for (int64_t i = 0; i < vectorType.getNumElements(); ++i) {
@@ -641,8 +640,10 @@ static Value scalarizeVectorOpHelper(Operation *op, ValueRange operands,
return LLVM::ExtractElementOp::create(rewriter, loc, operand, index);
};
auto scalarOperands = llvm::map_to_vector(operands, extractElement);
- Operation *scalarOp =
- rewriter.create(loc, name, scalarOperands, elementType, op->getAttrs());
+ OperationState state(loc, op->getName(), scalarOperands, elementType,
+ op->getDiscardableAttrDictionary().getValue());
+ state.propertiesAttr = op->getPropertiesAsAttribute();
+ Operation *scalarOp = rewriter.create(state);
result = LLVM::InsertElementOp::create(rewriter, loc, result,
scalarOp->getResult(0), index);
}
@@ -808,7 +809,8 @@ LogicalResult GPUReturnOpLowering::matchAndRewrite(
// If ReturnOp has 0 or 1 operand, create it and return immediately.
if (numArguments <= 1) {
rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(
- op, TypeRange(), updatedOperands, op->getAttrs());
+ op, TypeRange(), updatedOperands,
+ op->getDiscardableAttrDictionary().getValue());
return success();
}
@@ -824,8 +826,8 @@ LogicalResult GPUReturnOpLowering::matchAndRewrite(
for (auto [idx, operand] : llvm::enumerate(updatedOperands)) {
packed = LLVM::InsertValueOp::create(rewriter, loc, packed, operand, idx);
}
- rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(op, TypeRange(), packed,
- op->getAttrs());
+ rewriter.replaceOpWithNewOp<LLVM::ReturnOp>(
+ op, TypeRange(), packed, op->getDiscardableAttrDictionary().getValue());
return success();
}
diff --git a/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h b/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
index 777f3b92b941c..2a6036ab97bb6 100644
--- a/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
+++ b/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
@@ -87,7 +87,7 @@ struct OpLowering : public ConvertOpToLLVMPattern<Op> {
opBound = static_cast<uint32_t>(bound->getZExtValue());
if (auto range = getIndexOpRange(op, op.getDimension(), opBound, indexKind,
intrType, /*bitWidth=*/32))
- newOp->setAttr("range", range);
+ newOp->setInherentAttr(rewriter.getStringAttr("range"), range);
if (indexBitwidth > 32) {
newOp = LLVM::SExtOp::create(rewriter, loc,
diff --git a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
index 80420c26537c3..2b57e90efe7a5 100644
--- a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
+++ b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
@@ -513,8 +513,8 @@ struct LowerGpuOpsToNVVMOpsPass final
// Request C wrapper emission.
for (auto func : m.getOps<func::FuncOp>()) {
- func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
- UnitAttr::get(&getContext()));
+ func->setDiscardableAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
+ UnitAttr::get(&getContext()));
}
// Customize the bitwidth used for the device side index computations.
diff --git a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
index a3819df4f8a84..ab7666bc5d530 100644
--- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
+++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
@@ -302,7 +302,7 @@ struct GPUSubgroupIdOpToROCDL : ConvertOpToLLVMPattern<gpu::SubgroupIdOp> {
...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/218921
More information about the Mlir-commits
mailing list