[Mlir-commits] [mlir] [mlir][LLVM][GPU] Migrate to explicit split inherent/discardable attribute APIs access (PR #218921)
Mehdi Amini
llvmlistbot at llvm.org
Wed Aug 26 08:02:45 PDT 2026
https://github.com/joker-eph updated https://github.com/llvm/llvm-project/pull/218921
>From f01b38f808ee16951103bec952fc8bbd27487f97 Mon Sep 17 00:00:00 2001
From: Mehdi Amini <joker.eph at gmail.com>
Date: Thu, 20 Aug 2026 06:31:48 -0700
Subject: [PATCH] [mlir][LLVM][GPU] Migrate to explicit attribute APIs
Use discardable attribute APIs and typed operation accessors throughout the
LLVM and GPU dialect families, their conversions, translations, and tests.
Assisted-by: Codex
---
.../ArithCommon/AttrToLLVMConverter.h | 53 ++++++-------------
.../mlir/Conversion/LLVMCommon/Pattern.h | 5 +-
.../Conversion/LLVMCommon/VectorPattern.h | 7 ++-
mlir/include/mlir/Dialect/GPU/IR/GPUOps.td | 6 ++-
.../LLVMIR/BasicPtxBuilderInterface.td | 17 +++---
.../mlir/Dialect/LLVMIR/LLVMInterfaces.td | 11 ++++
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 3 +-
mlir/include/mlir/IR/SymbolInterfaces.td | 7 +--
.../AMDGPUToROCDL/AMDGPUToROCDL.cpp | 6 +--
.../Conversion/ArithToLLVM/ArithToLLVM.cpp | 8 +--
.../ComplexToLLVM/ComplexToLLVM.cpp | 4 +-
.../ControlFlowToLLVM/ControlFlowToLLVM.cpp | 16 ++++--
.../Conversion/GPUCommon/GPUOpsLowering.cpp | 14 ++---
.../GPUCommon/IndexIntrinsicsOpLowering.h | 2 +-
.../GPUToNVVM/LowerGpuOpsToNVVMOps.cpp | 4 +-
.../GPUToROCDL/LowerGpuOpsToROCDLOps.cpp | 11 ++--
.../LowerFunctionDiscardablesToLLVM.cpp | 3 +-
mlir/lib/Conversion/LLVMCommon/Pattern.cpp | 12 ++++-
mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp | 3 +-
.../Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp | 4 +-
.../Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp | 4 +-
mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp | 32 +++++++----
.../AMDGPU/Transforms/EmulateAtomics.cpp | 5 +-
.../AMDGPU/Transforms/MaskedloadToLoad.cpp | 4 +-
mlir/lib/Dialect/GPU/IR/GPUDialect.cpp | 44 ++++++++-------
.../GPU/IR/InferIntRangeInterfaceImpls.cpp | 4 +-
.../GPU/Transforms/EliminateBarriers.cpp | 2 +-
.../GPU/Transforms/KernelOutlining.cpp | 8 +--
.../GPU/Transforms/ParallelLoopMapper.cpp | 7 +--
mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp | 4 +-
mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp | 33 +++++++-----
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp | 27 ++++++----
.../Transforms/InlinerInterfaceImpl.cpp | 2 +-
.../LLVMIR/Transforms/RequestCWrappers.cpp | 5 +-
mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp | 5 +-
.../NVGPU/Transforms/MmaSyncTF32Transform.cpp | 2 +-
.../XeGPU/Transforms/XeGPUBlocking.cpp | 20 +++----
.../XeGPU/Transforms/XeGPULayoutImpl.cpp | 10 ++--
.../XeGPU/Transforms/XeGPUPropagateLayout.cpp | 10 ++--
.../Transforms/XeGPUSgToLaneDistribute.cpp | 13 +++--
.../Dialect/XeGPU/Transforms/XeGPUUnroll.cpp | 30 +++++------
.../Transforms/XeGPUWgToSgDistribute.cpp | 5 +-
mlir/lib/Dialect/XeGPU/Utils/XeGPUUtils.cpp | 27 +++++-----
mlir/lib/Target/LLVM/XeVM/Target.cpp | 3 +-
.../LLVMIR/LLVMIRToLLVMTranslation.cpp | 2 +-
.../LLVMIR/LLVMToLLVMIRTranslation.cpp | 2 +-
.../ROCDL/ROCDLToLLVMIRTranslation.cpp | 8 +--
mlir/lib/Target/LLVMIR/ModuleImport.cpp | 27 +++++-----
mlir/lib/Target/LLVMIR/ModuleTranslation.cpp | 12 +++--
.../LLVMIR/Transforms/TargetToDataLayout.cpp | 9 ++--
.../Transforms/TargetToTargetFeatures.cpp | 5 +-
mlir/python/mlir/dialects/gpu/__init__.py | 10 ++--
.../ArithToLLVM/attribute-storage.mlir | 32 +++++++++++
mlir/test/lib/Dialect/LLVM/TestPatterns.cpp | 5 +-
.../lib/Dialect/XeGPU/TestXeGPUTransforms.cpp | 5 +-
.../Target/LLVM/SerializeToLLVMBitcode.cpp | 5 +-
56 files changed, 370 insertions(+), 254 deletions(-)
create mode 100644 mlir/test/Conversion/ArithToLLVM/attribute-storage.mlir
diff --git a/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h b/mlir/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h
index feb74c86e349f..cdea5307aa72d 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,13 @@ 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()) {
+ srcOp->getName().populateInherentAttrs(srcOp, convertedAttr);
+ convertedAttr.erase(SourceOp::getFastMathAttrName());
+ auto arithFMFAttr = srcOp.getFastMathFlagsAttr();
if (arithFMFAttr) {
- StringRef targetAttrName = TargetOp::getFastmathAttrName();
- convertedAttr.set(targetAttrName,
+ convertedAttr.set(TargetOp::getFastmathAttrName(),
convertArithFastMathAttrToLLVM(arithFMFAttr));
}
}
@@ -78,17 +75,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 +108,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 +133,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..10e42706abd3f 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td
@@ -157,13 +157,16 @@ 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.
+ op->getName().walkInherentAttrs(
+ op, [&](llvm::StringRef, mlir::Attribute &attr) {
+ if (auto intAttr = dyn_cast<mlir::IntegerAttr>(attr)) {
+ ::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..c09edbc7cb317 100644
--- a/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
+++ b/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
@@ -133,6 +133,16 @@ static SmallVector<Value> flattenValues(ArrayRef<ValueRange> values) {
return result;
}
+/// Set attributes on an operation using its inherent/discardable split.
+static void setConvertedAttrs(Operation *op, DictionaryAttr attrs) {
+ for (NamedAttribute attr : attrs) {
+ if (op->getInherentAttr(attr.getName()).has_value())
+ op->setInherentAttr(attr.getName(), attr.getValue());
+ else
+ op->setDiscardableAttr(attr.getName(), attr.getValue());
+ }
+}
+
/// Convert the destination block signature (if necessary) and lower the branch
/// op to llvm.br.
struct BranchOpLowering : public ConvertOpToLLVMPattern<cf::BranchOp> {
@@ -148,12 +158,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);
+ setConvertedAttrs(newOp, attrs);
return success();
}
};
@@ -191,7 +201,7 @@ struct CondBranchOpLowering : public ConvertOpToLLVMPattern<cf::CondBranchOp> {
*convertedTrueBlock, *convertedFalseBlock);
// TODO: We should not just forward all attributes like that. But there are
// existing Flang tests that depend on this behavior.
- newOp->setDiscardableAttrs(attrs);
+ setConvertedAttrs(newOp, 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> {
op, dim, std::nullopt,
gpu::index_lowering::IndexKind::Block,
gpu::index_lowering::IntrType::Id, 32))
- tidOp->setAttr("range", range);
+ tidOp->setInherentAttr(rewriter.getStringAttr("range"), range);
};
setBoundFromContext(tidX, gpu::Dimension::x);
setBoundFromContext(tidY, gpu::Dimension::y);
@@ -737,16 +737,17 @@ struct LowerGpuOpsToROCDLOpsPass final
gpu::GPUModuleOp m = getOperation();
MLIRContext *ctx = m.getContext();
- auto llvmDataLayout = m->getAttrOfType<StringAttr>(
+ auto llvmDataLayout = m->getDiscardableAttrOfType<StringAttr>(
LLVM::LLVMDialect::getDataLayoutAttrName());
if (!llvmDataLayout) {
llvmDataLayout = StringAttr::get(ctx, amdgcnDataLayout);
- m->setAttr(LLVM::LLVMDialect::getDataLayoutAttrName(), llvmDataLayout);
+ m->setDiscardableAttr(LLVM::LLVMDialect::getDataLayoutAttrName(),
+ llvmDataLayout);
}
// Request C wrapper emission.
for (auto func : m.getOps<func::FuncOp>()) {
- func->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
- UnitAttr::get(ctx));
+ func->setDiscardableAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
+ UnitAttr::get(ctx));
}
FailureOr<amdgpu::Chipset> maybeChipset = amdgpu::Chipset::parse(chipset);
diff --git a/mlir/lib/Conversion/LLVMCommon/LowerFunctionDiscardablesToLLVM.cpp b/mlir/lib/Conversion/LLVMCommon/LowerFunctionDiscardablesToLLVM.cpp
index 5427877fdfd8d..859d96d7e2377 100644
--- a/mlir/lib/Conversion/LLVMCommon/LowerFunctionDiscardablesToLLVM.cpp
+++ b/mlir/lib/Conversion/LLVMCommon/LowerFunctionDiscardablesToLLVM.cpp
@@ -29,7 +29,8 @@ mlir::lowerDiscardableAttrsForLLVMFunc(FunctionOpInterface funcOp,
NamedAttrList inherentAttrs;
- for (const NamedAttribute &attr : funcOp->getDiscardableAttrs()) {
+ for (const NamedAttribute &attr :
+ funcOp->getDiscardableAttrDictionary().getValue()) {
StringRef attrName = attr.getName().strref();
if (odsAttrNames.contains(attrName)) {
diff --git a/mlir/lib/Conversion/LLVMCommon/Pattern.cpp b/mlir/lib/Conversion/LLVMCommon/Pattern.cpp
index 2e0d92c3ba847..32d67dd0680db 100644
--- a/mlir/lib/Conversion/LLVMCommon/Pattern.cpp
+++ b/mlir/lib/Conversion/LLVMCommon/Pattern.cpp
@@ -367,7 +367,17 @@ LogicalResult LLVM::detail::intrinsicRewrite(
auto callIntrOp = LLVM::CallIntrinsicOp::create(
rewriter, loc, resType, rewriter.getStringAttr(intrinsic), operands);
// Propagate attributes.
- callIntrOp->setAttrs(op->getAttrDictionary());
+ auto copyAttr = [&](StringAttr name, Attribute attr) {
+ if (callIntrOp->getInherentAttr(name).has_value())
+ callIntrOp->setInherentAttr(name, attr);
+ else
+ callIntrOp->setDiscardableAttr(name, attr);
+ };
+ for (NamedAttribute attr : op->getDiscardableAttrDictionary())
+ copyAttr(attr.getName(), attr.getValue());
+ op->getName().walkInherentAttrs(op, [&](StringRef name, Attribute &attr) {
+ copyAttr(rewriter.getStringAttr(name), attr);
+ });
if (numResults <= 1) {
// Directly replace the original op.
diff --git a/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp b/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
index 7143dcc19ff10..4b29df764be71 100644
--- a/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
+++ b/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
@@ -145,7 +145,8 @@ struct ConvertNativeFuncPattern final : public OpConversionPattern<Op> {
// convert arith fastmath attrs into attrs recognized by llvm.
arith::AttrConvertFastMathToLLVM<Op, LLVM::CallOp> fastAttrConverter(op);
mlir::NamedAttribute fastAttr = fastAttrConverter.getAttrs()[0];
- callOp->setAttr(fastAttr.getName(), fastAttr.getValue());
+ callOp.setFastmathFlagsAttr(
+ cast<LLVM::FastmathFlagsAttr>(fastAttr.getValue()));
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 b155505a2dae1..ae3a5f7e8e97a 100644
--- a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
+++ b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
@@ -353,7 +353,7 @@ struct MmaSyncOptoNVVM : public ConvertOpToLLVMPattern<nvgpu::MmaSyncOp> {
std::array<int64_t, 3> gemmShape = op.getMmaShapeAsArray();
// Tensor Cores (mma.sync) on F32 works only with TensorFloat32 (TF32).
- bool tf32Enabled = op->hasAttr(op.getTf32EnabledAttrName());
+ bool tf32Enabled = op.getTf32Enabled().value_or(false);
if (aType.getElementType().isF32() && !tf32Enabled)
return failure();
@@ -599,7 +599,7 @@ struct NVGPUMmaSparseSyncLowering
"could not infer the PTX type for the accumulator/result");
// Same as `mma.sync`, F32 works only with TensorFloat32 (TF32).
- bool tf32Enabled = op->hasAttr(op.getTf32EnabledAttrName());
+ bool tf32Enabled = op.getTf32Enabled().value_or(false);
if (aType.getElementType().isF32() && !tf32Enabled)
return failure();
diff --git a/mlir/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp b/mlir/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp
index f782be5502a8a..98edcc2bc429c 100644
--- a/mlir/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp
+++ b/mlir/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp
@@ -804,7 +804,7 @@ class LoadStoreToXeVMPattern : public OpConversionPattern<OpType> {
Value loaded =
LLVM::LoadOp::create(rewriter, loc, valOrResTy, basePtrLLVM);
// Set cache control attribute on the load operation.
- loaded.getDefiningOp()->setAttr(
+ loaded.getDefiningOp()->setDiscardableAttr(
"cache_control", xevm::LoadCacheControlAttr::get(
ctxt, translateLoadXeGPUCacheHint(
op.getL1Hint(), op.getL3Hint())));
@@ -832,7 +832,7 @@ class LoadStoreToXeVMPattern : public OpConversionPattern<OpType> {
auto storeOp =
LLVM::StoreOp::create(rewriter, loc, adaptor.getValue(), basePtrLLVM);
// Set cache control attribute on the store operation.
- storeOp.getOperation()->setAttr(
+ storeOp.getOperation()->setDiscardableAttr(
"cache_control", xevm::StoreCacheControlAttr::get(
ctxt, translateStoreXeGPUCacheHint(
op.getL1Hint(), op.getL3Hint())));
diff --git a/mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp b/mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp
index 27ae552e83670..3b350ba8b6857 100644
--- a/mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp
+++ b/mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp
@@ -257,8 +257,9 @@ static std::optional<StoreCacheControl> getCacheControl(BlockStoreOp op) {
}
static std::optional<LoadCacheControl> getCacheControl(LLVM::LoadOp op) {
- if (op->hasAttr("cache_control")) {
- auto attr = op->getAttrOfType<xevm::LoadCacheControlAttr>("cache_control");
+ if (op->hasDiscardableAttr("cache_control")) {
+ auto attr = op->getDiscardableAttrOfType<xevm::LoadCacheControlAttr>(
+ "cache_control");
if (!attr)
return std::nullopt;
return std::optional<LoadCacheControl>(attr.getValue());
@@ -267,8 +268,9 @@ static std::optional<LoadCacheControl> getCacheControl(LLVM::LoadOp op) {
}
static std::optional<StoreCacheControl> getCacheControl(LLVM::StoreOp op) {
- if (op->hasAttr("cache_control")) {
- auto attr = op->getAttrOfType<xevm::StoreCacheControlAttr>("cache_control");
+ if (op->hasDiscardableAttr("cache_control")) {
+ auto attr = op->getDiscardableAttrOfType<xevm::StoreCacheControlAttr>(
+ "cache_control");
if (!attr)
return std::nullopt;
return std::optional<StoreCacheControl>(attr.getValue());
@@ -549,7 +551,18 @@ static LLVM::CallOp createDeviceFunctionCall(
funcOp.setArgAttr(idx, attrName, rewriter.getUnitAttr());
auto callOp = LLVM::CallOp::create(rewriter, loc, funcOp, args);
- callOp->setAttrs(funcOp->getAttrs());
+ auto copyAttr = [&](StringAttr name, Attribute attr) {
+ if (callOp->getInherentAttr(name).has_value())
+ callOp->setInherentAttr(name, attr);
+ else
+ callOp->setDiscardableAttr(name, attr);
+ };
+ for (NamedAttribute attr : funcOp->getDiscardableAttrDictionary())
+ copyAttr(attr.getName(), attr.getValue());
+ funcOp->getName().walkInherentAttrs(
+ funcOp, [&](StringRef name, Attribute &attr) {
+ copyAttr(rewriter.getStringAttr(name), attr);
+ });
return callOp;
}
@@ -984,14 +997,15 @@ class LLVMLoadStoreToOCLPattern : public OpConversionPattern<OpType> {
LogicalResult
matchAndRewrite(OpType op, typename OpType::Adaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- if (!op->hasAttr("cache_control"))
+ if (!op->hasDiscardableAttr("cache_control"))
return failure();
auto *moduleOp = op->template getParentWithTrait<OpTrait::SymbolTable>();
std::optional<ArrayAttr> optCacheControls =
getCacheControlMetadata(rewriter, op);
if (!optCacheControls) {
- rewriter.modifyOpInPlace(op, [&]() { op->removeAttr("cache_control"); });
+ rewriter.modifyOpInPlace(
+ op, [&]() { op->removeDiscardableAttr("cache_control"); });
return success();
}
@@ -1007,7 +1021,7 @@ class LLVMLoadStoreToOCLPattern : public OpConversionPattern<OpType> {
// Replace the pointer operand with the annotated one.
rewriter.modifyOpInPlace(op, [&]() {
op->setOperand(ptrIdx, annotatedPtr);
- op->removeAttr("cache_control");
+ op->removeDiscardableAttr("cache_control");
});
return success();
}
@@ -1831,7 +1845,7 @@ void ::mlir::populateXeVMToLLVMConversionPatterns(ConversionTarget &target,
return addrSpace != 3;
}
// cache_control attribute should be converted.
- return !op->hasAttr("cache_control");
+ return !op->hasDiscardableAttr("cache_control");
});
target.addIllegalDialect<XeVMDialect>();
patterns.add<LoadStorePrefetchToOCLPattern<BlockLoad2dOp>,
diff --git a/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp b/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
index 332ac8cdb60c0..b59f7e97553c1 100644
--- a/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
+++ b/mlir/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
@@ -110,7 +110,10 @@ LogicalResult RawBufferAtomicByCasPattern<AtomicOp, ArithOp>::matchAndRewrite(
ConversionPatternRewriter &rewriter) const {
Location loc = atomicOp.getLoc();
- ArrayRef<NamedAttribute> origAttrs = atomicOp->getAttrs();
+ NamedAttrList origAttrs(atomicOp->getDiscardableAttrDictionary());
+ atomicOp->getName().walkInherentAttrs(
+ atomicOp,
+ [&](StringRef name, Attribute &attr) { origAttrs.append(name, attr); });
ValueRange operands = adaptor.getOperands();
Value data = operands.take_front()[0];
ValueRange invariantArgs = operands.drop_front();
diff --git a/mlir/lib/Dialect/AMDGPU/Transforms/MaskedloadToLoad.cpp b/mlir/lib/Dialect/AMDGPU/Transforms/MaskedloadToLoad.cpp
index 89ef51f922cad..8c2d697632e20 100644
--- a/mlir/lib/Dialect/AMDGPU/Transforms/MaskedloadToLoad.cpp
+++ b/mlir/lib/Dialect/AMDGPU/Transforms/MaskedloadToLoad.cpp
@@ -81,7 +81,7 @@ struct MaskedLoadLowering final : OpRewritePattern<vector::MaskedLoadOp> {
LogicalResult matchAndRewrite(vector::MaskedLoadOp maskedOp,
PatternRewriter &rewriter) const override {
- if (maskedOp->hasAttr(kMaskedloadNeedsMask))
+ if (maskedOp->hasDiscardableAttr(kMaskedloadNeedsMask))
return rewriter.notifyMatchFailure(maskedOp, "already rewritten");
if (failed(hasBufferAddressSpace(maskedOp.getBase().getType()))) {
@@ -150,7 +150,7 @@ struct MaskedLoadLowering final : OpRewritePattern<vector::MaskedLoadOp> {
auto thenBuilder = [&](OpBuilder &builder, Location loc) {
Operation *read = builder.clone(*maskedOp.getOperation());
- read->setAttr(kMaskedloadNeedsMask, builder.getUnitAttr());
+ read->setDiscardableAttr(kMaskedloadNeedsMask, builder.getUnitAttr());
Value readResult = read->getResult(0);
scf::YieldOp::create(builder, loc, readResult);
};
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index 22c5e3c9b86ad..4250f06c6a12d 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -266,7 +266,7 @@ bool GPUDialect::isKernel(Operation *op) {
if (auto gpuFunc = dyn_cast<GPUFuncOp>(op))
return gpuFunc.isKernel();
return static_cast<bool>(
- op->getAttrOfType<UnitAttr>(getKernelFuncAttrName()));
+ op->getDiscardableAttrOfType<UnitAttr>(getKernelFuncAttrName()));
}
namespace {
@@ -695,7 +695,8 @@ void gpu::addAsyncDependency(Operation *op, Value token) {
return;
auto attrName =
OpTrait::AttrSizedOperandSegments<void>::getOperandSegmentSizeAttr();
- auto sizeAttr = op->template getAttrOfType<DenseI32ArrayAttr>(attrName);
+ auto sizeAttr = dyn_cast_or_null<DenseI32ArrayAttr>(
+ op->getInherentAttr(attrName).value_or(Attribute{}));
// Async dependencies is the only variadic operand.
if (!sizeAttr)
@@ -703,7 +704,8 @@ void gpu::addAsyncDependency(Operation *op, Value token) {
SmallVector<int32_t, 8> sizes(sizeAttr.asArrayRef());
++sizes.front();
- op->setAttr(attrName, Builder(op->getContext()).getDenseI32ArrayAttr(sizes));
+ op->setInherentAttr(StringAttr::get(op->getContext(), attrName),
+ Builder(op->getContext()).getDenseI32ArrayAttr(sizes));
}
//===----------------------------------------------------------------------===//
@@ -961,11 +963,11 @@ void LaunchOp::print(OpAsmPrinter &p) {
p << ' ';
p.printRegion(getBody(), /*printEntryBlockArgs=*/false);
- p.printOptionalAttrDict((*this)->getAttrs(), /*elidedAttrs=*/{
- LaunchOp::getOperandSegmentSizeAttr(),
- getWorkgroupAttributionsAttrName(),
- getCooperativeAttrName(), moduleAttrName,
- functionAttrName});
+ p.printOptionalAttrDict(
+ (*this)->getDiscardableAttrDictionary().getValue(), /*elidedAttrs=*/{
+ LaunchOp::getOperandSegmentSizeAttr(),
+ getWorkgroupAttributionsAttrName(), getCooperativeAttrName(),
+ moduleAttrName, functionAttrName});
}
// Parse the size assignment blocks for blocks and threads. These have the form
@@ -1346,7 +1348,7 @@ LogicalResult LaunchFuncOp::verify() {
if (!module)
return emitOpError("expected to belong to a module");
- if (!module->getAttrOfType<UnitAttr>(
+ if (!module->getDiscardableAttrOfType<UnitAttr>(
GPUDialect::getContainerModuleAttrName()))
return emitOpError("expected the closest surrounding module to have the '" +
GPUDialect::getContainerModuleAttrName() +
@@ -1378,8 +1380,7 @@ LaunchFuncOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
// Ignore launch ops with missing attributes here. The errors will be
// reported by the verifiers of those ops.
- if (!launchOp->getAttrOfType<SymbolRefAttr>(
- LaunchFuncOp::getKernelAttrName(launchOp->getName())))
+ if (!launchOp.getKernelAttr())
return success();
// Check that `launch_func` refers to a well-formed GPU kernel container.
@@ -1813,7 +1814,9 @@ void GPUFuncOp::print(OpAsmPrinter &p) {
static DictionaryAttr getAttributionAttrs(GPUFuncOp op, unsigned index,
StringAttr attrName) {
- auto allAttrs = llvm::dyn_cast_or_null<ArrayAttr>(op->getAttr(attrName));
+ ArrayAttr allAttrs = attrName == op.getWorkgroupAttribAttrsAttrName()
+ ? op.getWorkgroupAttribAttrsAttr()
+ : op.getPrivateAttribAttrsAttr();
if (!allAttrs || index >= allAttrs.size())
return DictionaryAttr();
return llvm::cast<DictionaryAttr>(allAttrs[index]);
@@ -1830,7 +1833,9 @@ DictionaryAttr GPUFuncOp::getPrivateAttributionAttrs(unsigned index) {
static void setAttributionAttrs(GPUFuncOp op, unsigned index,
DictionaryAttr value, StringAttr attrName) {
MLIRContext *ctx = op.getContext();
- auto allAttrs = llvm::dyn_cast_or_null<ArrayAttr>(op->getAttr(attrName));
+ ArrayAttr allAttrs = attrName == op.getWorkgroupAttribAttrsAttrName()
+ ? op.getWorkgroupAttribAttrsAttr()
+ : op.getPrivateAttribAttrsAttr();
SmallVector<Attribute> elements;
if (allAttrs)
elements.append(allAttrs.begin(), allAttrs.end());
@@ -1841,7 +1846,10 @@ static void setAttributionAttrs(GPUFuncOp op, unsigned index,
else
elements[index] = value;
ArrayAttr newValue = ArrayAttr::get(ctx, elements);
- op->setAttr(attrName, newValue);
+ if (attrName == op.getWorkgroupAttribAttrsAttrName())
+ op.setWorkgroupAttribAttrsAttr(newValue);
+ else
+ op.setPrivateAttribAttrsAttr(newValue);
}
void GPUFuncOp::setworkgroupAttributionAttrs(unsigned index,
@@ -2027,7 +2035,7 @@ void GPUModuleOp::setTargets(ArrayRef<TargetAttrInterface> targets) {
}
LogicalResult GPUModuleOp::verify() {
- auto targets = getOperation()->getAttrOfType<ArrayAttr>("targets");
+ auto targets = getTargetsAttr();
if (!targets)
return success();
@@ -2453,8 +2461,7 @@ void WarpExecuteOnLane0Op::print(OpAsmPrinter &p) {
p << "(" << getLaneid() << ")";
SmallVector<StringRef> coreAttr = {getWarpSizeAttrName()};
- auto warpSizeAttr = getOperation()->getAttr(getWarpSizeAttrName());
- p << "[" << llvm::cast<IntegerAttr>(warpSizeAttr).getInt() << "]";
+ p << "[" << getWarpSize() << "]";
if (!getArgs().empty())
p << " args(" << getArgs() << " : " << getArgs().getTypes() << ")";
@@ -2464,7 +2471,8 @@ void WarpExecuteOnLane0Op::print(OpAsmPrinter &p) {
p.printRegion(getRegion(),
/*printEntryBlockArgs=*/true,
/*printBlockTerminators=*/!getResults().empty());
- p.printOptionalAttrDict(getOperation()->getAttrs(), coreAttr);
+ p.printOptionalAttrDict(
+ getOperation()->getDiscardableAttrDictionary().getValue(), coreAttr);
}
ParseResult WarpExecuteOnLane0Op::parse(OpAsmParser &parser,
diff --git a/mlir/lib/Dialect/GPU/IR/InferIntRangeInterfaceImpls.cpp b/mlir/lib/Dialect/GPU/IR/InferIntRangeInterfaceImpls.cpp
index 33695833b1ad8..b15deed706076 100644
--- a/mlir/lib/Dialect/GPU/IR/InferIntRangeInterfaceImpls.cpp
+++ b/mlir/lib/Dialect/GPU/IR/InferIntRangeInterfaceImpls.cpp
@@ -68,7 +68,9 @@ getKnownLaunchAttr(GPUFuncOp func, DimensionKind dims, Dimension dim) {
static std::optional<uint32_t> getKnownLaunchAttr(FunctionOpInterface func,
StringRef attrName,
Dimension dim) {
- auto bounds = func.getOperation()->getAttrOfType<DenseI32ArrayAttr>(attrName);
+ auto bounds =
+ func.getOperation()->getDiscardableAttrOfType<DenseI32ArrayAttr>(
+ attrName);
if (!bounds)
return std::nullopt;
if (bounds.size() <= static_cast<uint32_t>(dim))
diff --git a/mlir/lib/Dialect/GPU/Transforms/EliminateBarriers.cpp b/mlir/lib/Dialect/GPU/Transforms/EliminateBarriers.cpp
index 21162546fe1b2..05ede7b017031 100644
--- a/mlir/lib/Dialect/GPU/Transforms/EliminateBarriers.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/EliminateBarriers.cpp
@@ -45,7 +45,7 @@ using namespace mlir::gpu;
/// Returns `true` if the op is defines the parallel region that is subject to
/// barrier synchronization.
static bool isParallelRegionBoundary(Operation *op) {
- if (op->hasAttr("__parallel_region_boundary_for_test"))
+ if (op->hasDiscardableAttr("__parallel_region_boundary_for_test"))
return true;
return isa<GPUFuncOp, LaunchOp>(op);
diff --git a/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp b/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
index 831ac7fb0354e..0ee3d425f66be 100644
--- a/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
@@ -394,8 +394,9 @@ class GpuKernelOutliningPass
// If any new module was inserted in this module, annotate this module as
// a container module.
if (modified)
- getOperation()->setAttr(gpu::GPUDialect::getContainerModuleAttrName(),
- UnitAttr::get(&getContext()));
+ getOperation()->setDiscardableAttr(
+ gpu::GPUDialect::getContainerModuleAttrName(),
+ UnitAttr::get(&getContext()));
}
private:
@@ -429,7 +430,8 @@ class GpuKernelOutliningPass
// If a valid data layout spec was provided, attach it to the kernel module.
// Otherwise, the default data layout will be used.
if (dataLayoutSpec)
- kernelModule->setAttr(DLTIDialect::kDataLayoutAttrName, dataLayoutSpec);
+ kernelModule->setDiscardableAttr(DLTIDialect::kDataLayoutAttrName,
+ dataLayoutSpec);
SymbolTable symbolTable(kernelModule);
symbolTable.insert(kernelFunc);
diff --git a/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp b/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp
index 594c7a265667e..07b6dd482f172 100644
--- a/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp
@@ -44,8 +44,9 @@ gpu::setMappingAttr(ParallelOp ploopOp,
specifiedMappings.insert(processor);
}
ArrayRef<Attribute> mappingAsAttrs(mapping.data(), mapping.size());
- ploopOp->setAttr(getMappingAttrName(),
- ArrayAttr::get(ploopOp.getContext(), mappingAsAttrs));
+ ploopOp->setDiscardableAttr(
+ getMappingAttrName(),
+ ArrayAttr::get(ploopOp.getContext(), mappingAsAttrs));
return success();
}
@@ -126,7 +127,7 @@ static void
mapParallelOp(ParallelOp parallelOp, MappingLevel mappingLevel = MapGrid,
MappingPolicy mappingPolicy = MappingPolicy::OutermostFirst) {
// Do not try to add a mapping to already mapped loops or nested loops.
- if (parallelOp->getAttr(getMappingAttrName()) ||
+ if (parallelOp->getDiscardableAttr(getMappingAttrName()) ||
((mappingLevel == MapGrid) && parallelOp->getParentOfType<ParallelOp>()))
return;
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp
index b81b3121f2e18..fdfae67bce47a 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp
@@ -620,8 +620,8 @@ TargetFeaturesAttr TargetFeaturesAttr::featuresAt(Operation *op) {
auto parentFunction = op->getParentOfType<FunctionOpInterface>();
if (!parentFunction)
return {};
- return parentFunction.getOperation()->getAttrOfType<TargetFeaturesAttr>(
- getAttributeName());
+ return parentFunction.getOperation()
+ ->getDiscardableAttrOfType<TargetFeaturesAttr>(getAttributeName());
}
FailureOr<Attribute> TargetFeaturesAttr::query(DataLayoutEntryKey key) {
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
index 374ad4a9dcb83..86252a3aef584 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
@@ -50,6 +50,13 @@ using mlir::LLVM::tailcallkind::getMaxEnumValForTailCallKind;
static constexpr const char kElemTypeAttrName[] = "elem_type";
+static NamedAttrList getAttrsForPrinting(Operation *op) {
+ NamedAttrList attrs(op->getRawDictionaryAttrs());
+ op->getName().walkInherentAttrs(
+ op, [&](StringRef name, Attribute &attr) { attrs.set(name, attr); });
+ return NamedAttrList(attrs.getDictionary(op->getContext()));
+}
+
static auto processFMFAttr(ArrayRef<NamedAttribute> attrs) {
SmallVector<NamedAttribute, 8> filteredAttrs(
llvm::make_filter_range(attrs, [&](NamedAttribute attr) {
@@ -278,14 +285,15 @@ static std::optional<ParseResult> parseOpBundles(
void ICmpOp::print(OpAsmPrinter &p) {
p << " \"" << stringifyICmpPredicate(getPredicate()) << "\" " << getOperand(0)
<< ", " << getOperand(1);
- p.printOptionalAttrDict((*this)->getAttrs(), {"predicate"});
+ p.printOptionalAttrDict(getAttrsForPrinting(*this).getAttrs(), {"predicate"});
p << " : " << getLhs().getType();
}
void FCmpOp::print(OpAsmPrinter &p) {
p << " \"" << stringifyFCmpPredicate(getPredicate()) << "\" " << getOperand(0)
<< ", " << getOperand(1);
- p.printOptionalAttrDict(processFMFAttr((*this)->getAttrs()), {"predicate"});
+ p.printOptionalAttrDict(processFMFAttr(getAttrsForPrinting(*this).getAttrs()),
+ {"predicate"});
p << " : " << getLhs().getType();
}
@@ -397,13 +405,10 @@ void AllocaOp::print(OpAsmPrinter &p) {
p << " inalloca";
p << ' ' << getArraySize() << " x " << getElemType();
+ NamedAttrList attrs((*this)->getDiscardableAttrDictionary().getValue());
if (getAlignment() && *getAlignment() != 0)
- p.printOptionalAttrDict((*this)->getAttrs(),
- {kElemTypeAttrName, getInallocaAttrName()});
- else
- p.printOptionalAttrDict(
- (*this)->getAttrs(),
- {getAlignmentAttrName(), kElemTypeAttrName, getInallocaAttrName()});
+ attrs.append(getAlignmentAttrName(), getAlignmentAttr());
+ p.printOptionalAttrDict(attrs);
p << " : " << funcTy;
}
@@ -1379,7 +1384,7 @@ void CallOp::print(OpAsmPrinter &p) {
getOpBundleOperands().getTypes(), getOpBundleTags());
}
- p.printOptionalAttrDict(processFMFAttr((*this)->getAttrs()),
+ p.printOptionalAttrDict(processFMFAttr(getAttrsForPrinting(*this).getAttrs()),
{getCalleeAttrName(), getTailCallKindAttrName(),
getVarCalleeTypeAttrName(), getCConvAttrName(),
getOperandSegmentSizesAttrName(),
@@ -1707,7 +1712,7 @@ void InvokeOp::print(OpAsmPrinter &p) {
getOpBundleOperands().getTypes(), getOpBundleTags());
}
- p.printOptionalAttrDict((*this)->getAttrs(),
+ p.printOptionalAttrDict(getAttrsForPrinting(*this).getAttrs(),
{getCalleeAttrName(), getOperandSegmentSizeAttr(),
getCConvAttrName(), getVarCalleeTypeAttrName(),
getOpBundleSizesAttrName(),
@@ -1887,7 +1892,7 @@ void LandingpadOp::print(OpAsmPrinter &p) {
<< value.getType() << ") ";
}
- p.printOptionalAttrDict((*this)->getAttrs(), {"cleanup"});
+ p.printOptionalAttrDict(getAttrsForPrinting(*this).getAttrs(), {"cleanup"});
p << ": " << getType();
}
@@ -2480,7 +2485,7 @@ void GlobalOp::print(OpAsmPrinter &p) {
// Note that the alignment attribute is printed using the
// default syntax here, even though it is an inherent attribute
// (as defined in https://mlir.llvm.org/docs/LangRef/#attributes)
- p.printOptionalAttrDict((*this)->getAttrs(),
+ p.printOptionalAttrDict(getAttrsForPrinting(*this).getAttrs(),
{SymbolTable::getSymbolAttrName(),
getGlobalTypeAttrName(), getConstantAttrName(),
getValueAttrName(), getLinkageAttrName(),
@@ -2841,7 +2846,7 @@ void AliasOp::print(OpAsmPrinter &p) {
printCommonGlobalAndAlias<AliasOp>(p, *this);
p.printSymbolName(getSymName());
- p.printOptionalAttrDict((*this)->getAttrs(),
+ p.printOptionalAttrDict(getAttrsForPrinting(*this).getAttrs(),
{SymbolTable::getSymbolAttrName(),
getAliasTypeAttrName(), getLinkageAttrName(),
getUnnamedAddrAttrName(), getTlsModeAttrName(),
@@ -4204,7 +4209,7 @@ void CallIntrinsicOp::print(OpAsmPrinter &p) {
getOpBundleOperands().getTypes(), getOpBundleTagsAttr());
}
- p.printOptionalAttrDict(processFMFAttr((*this)->getAttrs()),
+ p.printOptionalAttrDict(processFMFAttr(getAttrsForPrinting(*this).getAttrs()),
{getOperandSegmentSizesAttrName(),
getOpBundleSizesAttrName(), getIntrinAttrName(),
getOpBundleTagsAttrName(), getArgAttrsAttrName(),
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 4dddb33fd050e..dff60a22f48a0 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -1009,7 +1009,8 @@ void MmaOp::print(OpAsmPrinter &p) {
getLayoutBAttrName(),
getMultiplicandAPtxTypeAttrName(),
getMultiplicandBPtxTypeAttrName()});
- p.printOptionalAttrDict(this->getOperation()->getAttrs(), ignoreAttrNames);
+ p.printOptionalAttrDict((*this)->getDiscardableAttrDictionary().getValue(),
+ ignoreAttrNames);
// Print the types of the operands and result.
p << " : " << "(";
@@ -1502,7 +1503,8 @@ void MmaSpOp::print(OpAsmPrinter &p) {
getMultiplicandAPtxTypeAttrName(),
getMultiplicandBPtxTypeAttrName(),
getOrderedMetadataAttrName(), getKindAttrName()});
- p.printOptionalAttrDict((*this)->getAttrs(), ignoreAttrNames);
+ p.printOptionalAttrDict((*this)->getDiscardableAttrDictionary().getValue(),
+ ignoreAttrNames);
p << " : ";
p << "(";
for (int i = 0; i < 3; ++i) {
@@ -2092,7 +2094,8 @@ void MmaBlockScaleOp::print(OpAsmPrinter &p) {
getMultiplicandBPtxTypeAttrName(),
getScaleVecSizeAttrName(),
getBlockScaleFormatAttrName(), getKindAttrName()});
- p.printOptionalAttrDict(this->getOperation()->getAttrs(), ignoreAttrNames);
+ p.printOptionalAttrDict((*this)->getDiscardableAttrDictionary().getValue(),
+ ignoreAttrNames);
// Print type signature
p << " : (";
@@ -2350,7 +2353,8 @@ void MmaSpBlockScaleOp::print(OpAsmPrinter &p) {
getOrderedMetadataAttrName(),
getScaleVecSizeAttrName(),
getBlockScaleFormatAttrName(), getKindAttrName()});
- p.printOptionalAttrDict(this->getOperation()->getAttrs(), ignoreAttrNames);
+ p.printOptionalAttrDict((*this)->getDiscardableAttrDictionary().getValue(),
+ ignoreAttrNames);
// Print type signature
p << " : (";
@@ -5600,12 +5604,13 @@ LogicalResult Tcgen05StOp::verify() {
/// Infer the result ranges for the NVVM SpecialRangeableRegisterOp that might
/// have ConstantRangeAttr.
-static void nvvmInferResultRanges(Operation *op, Value result,
+static void nvvmInferResultRanges(std::optional<LLVM::ConstantRangeAttr> range,
+ Value result,
ArrayRef<::mlir::ConstantIntRanges> argRanges,
SetIntRangeFn setResultRanges) {
- if (auto rangeAttr = op->getAttrOfType<LLVM::ConstantRangeAttr>("range")) {
- setResultRanges(result, {rangeAttr.getLower(), rangeAttr.getUpper(),
- rangeAttr.getLower(), rangeAttr.getUpper()});
+ if (range) {
+ setResultRanges(result, {range->getLower(), range->getUpper(),
+ range->getLower(), range->getUpper()});
} else {
setResultRanges(result, IntegerValueRange::getMaxRange(result).getValue());
}
@@ -6666,8 +6671,8 @@ LogicalResult NVVMDialect::verifyOperationAttribute(Operation *op,
}
// blocksareclusters must be used along with reqntid and cluster_dim
if (attrName == NVVMDialect::getBlocksAreClustersAttrName()) {
- if (!op->hasAttr(NVVMDialect::getReqntidAttrName()) ||
- !op->hasAttr(NVVMDialect::getClusterDimAttrName())) {
+ if (!op->hasDiscardableAttr(NVVMDialect::getReqntidAttrName()) ||
+ !op->hasDiscardableAttr(NVVMDialect::getClusterDimAttrName())) {
return op->emitError()
<< "'" << attrName << "' attribute must be used along with " << "'"
<< NVVMDialect::getReqntidAttrName() << "' and " << "'"
@@ -6686,7 +6691,7 @@ LogicalResult NVVMDialect::verifyRegionArgAttribute(Operation *op,
if (!funcOp)
return success();
- bool isKernel = op->hasAttr(NVVMDialect::getKernelFuncAttrName());
+ bool isKernel = op->hasDiscardableAttr(NVVMDialect::getKernelFuncAttrName());
StringAttr attrName = argAttr.getName();
if (attrName == NVVM::NVVMDialect::getGridConstantAttrName()) {
if (!isKernel) {
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
index 657835ae54149..9409141f6e45a 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
@@ -286,7 +286,7 @@ static void createNewAliasScopesFromNoAliasParameter(
continue;
ssaCopies.insert(ssaCopy);
- if (!ssaCopy->hasAttr(LLVM::LLVMDialect::getNoAliasAttrName()))
+ if (!ssaCopy->hasDiscardableAttr(LLVM::LLVMDialect::getNoAliasAttrName()))
continue;
noAliasParams.insert(ssaCopy);
}
diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp
index a5fccbdc368fe..10d203c470f1a 100644
--- a/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp
+++ b/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp
@@ -25,8 +25,9 @@ class RequestCWrappersPass
: public LLVM::impl::LLVMRequestCWrappersPassBase<RequestCWrappersPass> {
public:
void runOnOperation() override {
- getOperation()->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
- UnitAttr::get(&getContext()));
+ getOperation()->setDiscardableAttr(
+ LLVM::LLVMDialect::getEmitCWrapperAttrName(),
+ UnitAttr::get(&getContext()));
}
};
} // namespace
diff --git a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
index 642e588488641..2649564a7828e 100644
--- a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
+++ b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
@@ -272,7 +272,7 @@ LogicalResult MmaSyncOp::verify() {
return verifyMmaSyncOp(this->getOperation(), getMatrixA(), getMatrixB(),
getMatrixC(), getMmaShapeAsArray(),
- getOperation()->hasAttr(getTf32EnabledAttrName()));
+ getTf32Enabled().value_or(false));
}
//===----------------------------------------------------------------------===//
@@ -296,8 +296,7 @@ LogicalResult MmaSparseSyncOp::verify() {
return verifyMmaSyncOp(this->getOperation(), getMatrixA(), getMatrixB(),
getMatrixC(), getMmaShapeAsArray(),
- getOperation()->hasAttr(getTf32EnabledAttrName()),
- true);
+ getTf32Enabled().value_or(false), true);
}
//===----------------------------------------------------------------------===//
diff --git a/mlir/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp b/mlir/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp
index 7f626a625aaea..118796e23a69a 100644
--- a/mlir/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp
+++ b/mlir/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp
@@ -35,7 +35,7 @@ struct MmaSyncF32ToTF32Pattern : public OpRewritePattern<nvgpu::MmaSyncOp> {
PatternRewriter &rewriter) const override {
Location location = op->getLoc();
- if (op->hasAttr(op.getTf32EnabledAttrName()) ||
+ if (op.getTf32Enabled().value_or(false) ||
!cast<VectorType>(op.getMatrixA().getType()).getElementType().isF32())
return failure();
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp
index 9d1686403a431..f70ff26b3d8f7 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp
@@ -433,14 +433,15 @@ void XeGPUBlockingPass::runOnOperation() {
if (!isa<scf::ForOp, scf::WhileOp, scf::ConditionOp, scf::IfOp>(loopOp))
return;
SmallVector<StringRef> toRemove;
- for (const NamedAttribute &attr : loopOp->getAttrs()) {
+ for (const NamedAttribute &attr :
+ loopOp->getDiscardableAttrDictionary().getValue()) {
StringRef name = attr.getName().strref();
if (name.starts_with("layout_operand_") ||
name.starts_with("layout_result_"))
toRemove.push_back(name);
}
for (StringRef name : toRemove)
- loopOp->removeAttr(name);
+ loopOp->removeDiscardableAttr(name);
});
// Source (N:1) and target (1:1) materializations using
@@ -524,15 +525,16 @@ void XeGPUBlockingPass::runOnOperation() {
// Remove the layout attributes cached per operands.
for (OpOperand &opr : op->getOpOperands()) {
std::string name = xegpu::getTemporaryLayoutName(opr);
- if (op->hasAttrOfType<xegpu::DistributeLayoutAttr>(name))
- op->removeAttr(name);
+ if (op->hasDiscardableAttrOfType<xegpu::DistributeLayoutAttr>(name))
+ op->removeDiscardableAttr(name);
}
// Update the layout attributes per result.
for (OpResult result : op->getOpResults()) {
std::string name = xegpu::getTemporaryLayoutName(result);
- if (auto layout = op->getAttrOfType<xegpu::DistributeLayoutAttr>(name)) {
- op->removeAttr(name);
+ if (auto layout =
+ op->getDiscardableAttrOfType<xegpu::DistributeLayoutAttr>(name)) {
+ op->removeDiscardableAttr(name);
if (!isa<LoopLikeOpInterface>(op))
xegpu::setDistributeLayoutAttr(result, layout.dropInstData());
}
@@ -540,9 +542,9 @@ void XeGPUBlockingPass::runOnOperation() {
// Drop left-over inst_data if the unroll pattern does not being applied,
// say, inst_data just matches their shape.
- SmallVector<NamedAttribute> newAttrs =
- xegpu::dropInstDataOnAttrs(op->getAttrs());
- op->setAttrs(newAttrs);
+ SmallVector<NamedAttribute> newAttrs = xegpu::dropInstDataOnAttrs(
+ op->getDiscardableAttrDictionary().getValue());
+ op->setDiscardableAttrs(newAttrs);
});
// Resolve UnrealizedConversionCastOps generated by SCF structural type
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp
index b11f7ecd3df06..4825d12d347c4 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp
@@ -377,8 +377,8 @@ template <typename T, typename>
void xegpu::removeLayoutAttr(const T &operandOrResult) {
Operation *owner = operandOrResult.getOwner();
std::string name = xegpu::getTemporaryLayoutName(operandOrResult);
- if (owner->hasAttrOfType<DistributeLayoutAttr>(name))
- owner->removeAttr(name);
+ if (owner->hasDiscardableAttrOfType<DistributeLayoutAttr>(name))
+ owner->removeDiscardableAttr(name);
}
// Explicit instantiation for OpResult
@@ -393,19 +393,19 @@ void xegpu::removeLayoutAttrs(Operation *op) {
op->walk([&](Operation *nestOp) {
// Remove all attributes of DistributeLayoutAttr type
SmallVector<StringAttr> attrsToRemove;
- for (auto namedAttr : nestOp->getAttrs()) {
+ for (auto namedAttr : nestOp->getDiscardableAttrDictionary().getValue()) {
if (isa<DistributeLayoutAttr>(namedAttr.getValue()))
attrsToRemove.push_back(namedAttr.getName());
}
for (auto attrName : attrsToRemove)
- nestOp->removeAttr(attrName);
+ nestOp->removeDiscardableAttr(attrName);
});
}
void xegpu::removeTemporaryLayoutAttrs(Operation *op) {
op->walk([&](Operation *nestOp) {
SmallVector<StringAttr> attrsToRemove;
- for (auto namedAttr : nestOp->getDiscardableAttrs()) {
+ for (auto namedAttr : nestOp->getDiscardableAttrDictionary().getValue()) {
if (isa<xegpu::DistributeLayoutAttr>(namedAttr.getValue()))
attrsToRemove.push_back(namedAttr.getName());
}
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp
index 19d02e02b02ba..b7f44f07804c3 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp
@@ -1753,10 +1753,12 @@ ResolveLayoutConflicts::resolveTensorDescConsumer(OpOperand &operand) {
conflictingCreateNdOp.getContext(), currTDescType.getShape(),
currTDescType.getElementType(), currTDescType.getEncoding(),
expectedLayout);
- xegpu::CreateNdDescOp newOp = xegpu::CreateNdDescOp::create(
- builder, consumerOp->getLoc(), newTensorDescType,
- conflictingCreateNdOp->getOperands(),
- conflictingCreateNdOp->getAttrs());
+ OperationState state(
+ consumerOp->getLoc(), conflictingCreateNdOp->getName(),
+ conflictingCreateNdOp->getOperands(), TypeRange{newTensorDescType},
+ conflictingCreateNdOp->getDiscardableAttrDictionary().getValue());
+ state.propertiesAttr = conflictingCreateNdOp->getPropertiesAsAttribute();
+ auto newOp = cast<xegpu::CreateNdDescOp>(builder.create(state));
// Replace the tensor descriptor operand in the consumer op with the new
// tensor descriptor.
consumerOp->replaceUsesOfWith(tdescValue, newOp.getResult());
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp
index bc85140f9f121..b84fe9aa97b5c 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp
@@ -127,10 +127,12 @@ struct SgToLaneCreateNdDesc
if (!resultType.getLayout())
return failure();
- auto newOp = xegpu::CreateNdDescOp::create(
- rewriter, op.getLoc(), resultType.dropLayouts(), op.getOperands(),
- op->getAttrs());
- rewriter.replaceOp(op, newOp.getResult());
+ OperationState state(op.getLoc(), op->getName(), op.getOperands(),
+ TypeRange{resultType.dropLayouts()},
+ op->getDiscardableAttrDictionary().getValue());
+ state.propertiesAttr = op->getPropertiesAsAttribute();
+ Operation *newOp = rewriter.create(state);
+ rewriter.replaceOp(op, newOp->getResult(0));
return success();
}
};
@@ -345,10 +347,11 @@ struct SgToLaneElementWise : public ConversionPattern {
state.addOperands(operands);
state.addTypes(newResultType);
// Copy all attributes except for DistributeLayoutAttr.
- for (auto attr : op->getAttrs()) {
+ for (auto attr : op->getDiscardableAttrDictionary().getValue()) {
if (!isa<xegpu::DistributeLayoutAttr>(attr.getValue()))
state.addAttribute(attr.getName(), attr.getValue());
}
+ state.propertiesAttr = op->getPropertiesAsAttribute();
Operation *newOp = rewriter.create(state);
rewriter.replaceOp(op, newOp->getResult(0));
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp
index 5914e23f1f11a..2be5d21e5f46d 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp
@@ -428,9 +428,12 @@ struct UnrollDpasOp : public UnrollPattern<xegpu::DpasOp> {
if (tmpC)
operands.push_back(tmpC);
- tmpC = xegpu::DpasOp::create(
- rewriter, loc, vecTy, operands,
- xegpu::dropInstDataOnAttrs(op->getAttrs()));
+ OperationState state(
+ loc, op->getName(), operands, TypeRange{vecTy},
+ xegpu::dropInstDataOnAttrs(
+ op->getDiscardableAttrDictionary().getValue()));
+ state.propertiesAttr = op->getPropertiesAsAttribute();
+ tmpC = rewriter.create(state)->getResult(0);
}
newOps.push_back(tmpC);
}
@@ -530,9 +533,12 @@ struct UnrollDpasMxOp : public UnrollPattern<xegpu::DpasMxOp> {
operands.push_back(
bScaleVals[batch * (kIters * nIters) + k * nIters + j]);
- newDpasMxOp = xegpu::DpasMxOp::create(
- rewriter, loc, vecTy, operands,
- xegpu::dropInstDataOnAttrs(op->getAttrs()));
+ OperationState state(
+ loc, op->getName(), operands, TypeRange{vecTy},
+ xegpu::dropInstDataOnAttrs(
+ op->getDiscardableAttrDictionary().getValue()));
+ state.propertiesAttr = op->getPropertiesAsAttribute();
+ newDpasMxOp = cast<xegpu::DpasMxOp>(rewriter.create(state));
tmpC = newDpasMxOp.getResult();
}
newOps.push_back(newDpasMxOp);
@@ -563,11 +569,7 @@ struct UnrollLoadGatherOp : public UnrollPattern<xegpu::LoadGatherOp> {
return failure();
SmallVector<int64_t> targetMaskShape(*targetShape);
- int64_t chunkSize = 1;
- if (auto chunkSizeAttr = op->getAttr("chunk_size")) {
- if (auto intAttr = llvm::dyn_cast<IntegerAttr>(chunkSizeAttr))
- chunkSize = intAttr.getInt();
- }
+ int64_t chunkSize = op.getChunkSize().value_or(1);
// Unroll mask and offsets with correct shape
VectorType maskTy = llvm::dyn_cast<VectorType>(mask.getType());
@@ -657,11 +659,7 @@ struct UnrollStoreScatterOp : public UnrollPattern<xegpu::StoreScatterOp> {
if (!targetShape)
return failure();
- int64_t chunkSize = 1;
- if (auto chunkSizeAttr = op->getAttr("chunk_size")) {
- if (auto intAttr = llvm::dyn_cast<IntegerAttr>(chunkSizeAttr))
- chunkSize = intAttr.getInt();
- }
+ int64_t chunkSize = op.getChunkSize().value_or(1);
SmallVector<int64_t> targetMaskShape(*targetShape);
VectorType maskTy = llvm::dyn_cast<VectorType>(mask.getType());
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
index 0e8a386fb08b6..accc401b573bd 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
@@ -41,7 +41,7 @@ static xegpu::RangeAttr getRangeSpecAttr(Operation *op) {
Operation *parent = op->getParentOfType<scf::IfOp>();
while (parent) {
if (auto attr = llvm::dyn_cast_if_present<xegpu::RangeAttr>(
- parent->getAttr("sg_id_range")))
+ parent->getDiscardableAttr("sg_id_range")))
return attr;
parent = parent->getParentOfType<scf::IfOp>();
}
@@ -489,7 +489,8 @@ struct WgToSgElementwiseOp : public ConversionPattern {
OperationState state(op->getLoc(), op->getName());
state.addOperands(opOperands);
state.addTypes(newResultType);
- state.addAttributes(op->getAttrs());
+ state.addAttributes(op->getDiscardableAttrDictionary().getValue());
+ state.propertiesAttr = op->getPropertiesAsAttribute();
Operation *newOp = rewriter.create(state);
xegpu::removeLayoutAttrs(newOp);
newResults.push_back(newOp->getResult(0));
diff --git a/mlir/lib/Dialect/XeGPU/Utils/XeGPUUtils.cpp b/mlir/lib/Dialect/XeGPU/Utils/XeGPUUtils.cpp
index 76269cf193d13..734d48bebd8e0 100644
--- a/mlir/lib/Dialect/XeGPU/Utils/XeGPUUtils.cpp
+++ b/mlir/lib/Dialect/XeGPU/Utils/XeGPUUtils.cpp
@@ -160,9 +160,10 @@ xegpu::DistributeLayoutAttr xegpu::getDistributeLayoutAttr(const Value value) {
}
std::string layoutName = getTemporaryLayoutName(result);
- if (defOp->hasAttr(layoutName)) {
+ if (defOp->hasDiscardableAttr(layoutName)) {
auto layout =
- defOp->getAttrOfType<xegpu::DistributeLayoutAttr>(layoutName);
+ defOp->getDiscardableAttrOfType<xegpu::DistributeLayoutAttr>(
+ layoutName);
return layout;
}
}
@@ -254,8 +255,9 @@ xegpu::getDistributeLayoutAttr(const OpOperand &opr) {
}
std::string layoutName = xegpu::getTemporaryLayoutName(opr);
- if (op->hasAttr(layoutName)) {
- auto layout = op->getAttrOfType<xegpu::DistributeLayoutAttr>(layoutName);
+ if (op->hasDiscardableAttr(layoutName)) {
+ auto layout =
+ op->getDiscardableAttrOfType<xegpu::DistributeLayoutAttr>(layoutName);
return layout;
}
@@ -312,11 +314,11 @@ void xegpu::setDistributeLayoutAttr(
}
std::string name = xegpu::getTemporaryLayoutName(result);
- if (owner->hasAttrOfType<DistributeLayoutAttr>(name)) {
+ if (owner->hasDiscardableAttrOfType<DistributeLayoutAttr>(name)) {
return;
}
if (layout) {
- owner->setAttr(name, layout);
+ owner->setDiscardableAttr(name, layout);
}
}
@@ -360,11 +362,11 @@ void xegpu::setDistributeLayoutAttr(const OpOperand &operand,
}
std::string name = xegpu::getTemporaryLayoutName(operand);
- if (owner->hasAttrOfType<DistributeLayoutAttr>(name)) {
+ if (owner->hasDiscardableAttrOfType<DistributeLayoutAttr>(name)) {
return;
}
if (layout) {
- owner->setAttr(name, layout);
+ owner->setDiscardableAttr(name, layout);
}
}
@@ -374,8 +376,9 @@ xegpu::getTemporaryLayout(const T &operandOrResult) {
Operation *op = operandOrResult.getOwner();
std::string layoutName = xegpu::getTemporaryLayoutName(operandOrResult);
- if (op->hasAttr(layoutName)) {
- auto layout = op->getAttrOfType<xegpu::DistributeLayoutAttr>(layoutName);
+ if (op->hasDiscardableAttr(layoutName)) {
+ auto layout =
+ op->getDiscardableAttrOfType<xegpu::DistributeLayoutAttr>(layoutName);
return layout;
}
@@ -392,11 +395,11 @@ void xegpu::setTemporaryLayout(const T &operandOrResult,
const xegpu::DistributeLayoutAttr layout) {
Operation *owner = operandOrResult.getOwner();
std::string name = xegpu::getTemporaryLayoutName(operandOrResult);
- if (owner->hasAttrOfType<xegpu::DistributeLayoutAttr>(name)) {
+ if (owner->hasDiscardableAttrOfType<xegpu::DistributeLayoutAttr>(name)) {
return;
}
if (layout) {
- owner->setAttr(name, layout);
+ owner->setDiscardableAttr(name, layout);
}
}
diff --git a/mlir/lib/Target/LLVM/XeVM/Target.cpp b/mlir/lib/Target/LLVM/XeVM/Target.cpp
index 18f6b7b3e67c8..731ac4bd677f9 100644
--- a/mlir/lib/Target/LLVM/XeVM/Target.cpp
+++ b/mlir/lib/Target/LLVM/XeVM/Target.cpp
@@ -619,7 +619,8 @@ XeVMTargetAttrImpl::serializeToObject(Attribute attribute, Operation *module,
auto xeTarget = cast<XeVMTargetAttr>(attribute);
if (xeTarget.getTriple().starts_with("spirv")) {
gpuMod.walk([&](LLVM::LLVMFuncOp funcOp) {
- if (funcOp->hasAttr(gpu::GPUDialect::getKernelFuncAttrName())) {
+ if (funcOp->hasDiscardableAttr(
+ gpu::GPUDialect::getKernelFuncAttrName())) {
funcOp.setIntelReqdSubGroupSize(16);
return WalkResult::interrupt();
}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMIRToLLVMTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMIRToLLVMTranslation.cpp
index b6dee4e89ee4c..6733d6ee3ceb0 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMIRToLLVMTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMIRToLLVMTranslation.cpp
@@ -285,7 +285,7 @@ static LogicalResult setMmraAttr(llvm::MDNode *node, Operation *op,
}
mlirMmra = ArrayAttr::get(ctx, tags);
}
- op->setAttr(LLVMDialect::getMmraAttrName(), mlirMmra);
+ op->setDiscardableAttr(LLVMDialect::getMmraAttrName(), mlirMmra);
return success();
}
diff --git a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
index 3e0a228c773d5..aa62d0d0db4b1 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
@@ -639,7 +639,7 @@ convertOperationImpl(Operation &opInst, llvm::IRBuilderBase &builder,
invOp.getOpBundleTags(), moduleTranslation);
ArrayRef<llvm::Value *> operandsRef(operands);
llvm::InvokeInst *result;
- if (auto attr = opInst.getAttrOfType<FlatSymbolRefAttr>("callee")) {
+ if (auto attr = invOp.getCalleeAttr()) {
if (llvm::Function *function =
moduleTranslation.lookupFunction(attr.getValue())) {
result = builder.CreateInvoke(
diff --git a/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp b/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
index 42393218f8705..18eb07a9e17df 100644
--- a/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
@@ -195,8 +195,8 @@ class ROCDLDialectLLVMIRTranslationInterface
StringRef flatAttrName =
dialect->getFlatWorkGroupSizeAttrHelper().getName();
- if (auto flatAttr =
- dyn_cast_if_present<StringAttr>(op->getAttr(flatAttrName))) {
+ if (auto flatAttr = dyn_cast_if_present<StringAttr>(
+ op->getDiscardableAttr(flatAttrName))) {
if (flatAttr.getValue() != expectedFlatWorkGroupSize)
return op->emitOpError(Twine(flatAttrName) +
" must match rocdl.reqd_work_group_size");
@@ -204,8 +204,8 @@ class ROCDLDialectLLVMIRTranslationInterface
StringRef maxFlatAttrName =
dialect->getMaxFlatWorkGroupSizeAttrHelper().getName();
- if (auto maxFlatAttr =
- dyn_cast_if_present<IntegerAttr>(op->getAttr(maxFlatAttrName))) {
+ if (auto maxFlatAttr = dyn_cast_if_present<IntegerAttr>(
+ op->getDiscardableAttr(maxFlatAttrName))) {
llvm::SmallString<16> expectedMaxFlatWorkGroupSize;
llvm::raw_svector_ostream maxAttrValueStream(
expectedMaxFlatWorkGroupSize);
diff --git a/mlir/lib/Target/LLVMIR/ModuleImport.cpp b/mlir/lib/Target/LLVMIR/ModuleImport.cpp
index 08b23c0bbdfa4..e16125edb65d4 100644
--- a/mlir/lib/Target/LLVMIR/ModuleImport.cpp
+++ b/mlir/lib/Target/LLVMIR/ModuleImport.cpp
@@ -1025,8 +1025,9 @@ LogicalResult ModuleImport::convertDependentLibrariesMetadata() {
libraries.push_back(mdString->getString());
}
if (!libraries.empty())
- mlirModule->setAttr(LLVM::LLVMDialect::getDependentLibrariesAttrName(),
- builder.getStrArrayAttr(libraries));
+ mlirModule->setDiscardableAttr(
+ LLVM::LLVMDialect::getDependentLibrariesAttrName(),
+ builder.getStrArrayAttr(libraries));
}
return success();
}
@@ -1042,8 +1043,9 @@ LogicalResult ModuleImport::convertIdentMetadata() {
if (auto *md = dyn_cast<llvm::MDNode>(named.getOperand(0)))
if (md->getNumOperands() == 1)
if (auto *mdStr = dyn_cast<llvm::MDString>(md->getOperand(0)))
- mlirModule->setAttr(LLVMDialect::getIdentAttrName(),
- builder.getStringAttr(mdStr->getString()));
+ mlirModule->setDiscardableAttr(
+ LLVMDialect::getIdentAttrName(),
+ builder.getStringAttr(mdStr->getString()));
}
return success();
}
@@ -1059,8 +1061,9 @@ LogicalResult ModuleImport::convertCommandlineMetadata() {
if (auto *md = dyn_cast<llvm::MDNode>(nmd.getOperand(0)))
if (md->getNumOperands() == 1)
if (auto *mdStr = dyn_cast<llvm::MDString>(md->getOperand(0)))
- mlirModule->setAttr(LLVMDialect::getCommandlineAttrName(),
- builder.getStringAttr(mdStr->getString()));
+ mlirModule->setDiscardableAttr(
+ LLVMDialect::getCommandlineAttrName(),
+ builder.getStringAttr(mdStr->getString()));
}
return success();
}
@@ -1179,13 +1182,13 @@ LogicalResult ModuleImport::convertDataLayout() {
for (StringRef token : dataLayoutImporter.getUnhandledTokens())
emitWarning(loc, "unhandled data layout token: ") << token;
- mlirModule->setAttr(DLTIDialect::kDataLayoutAttrName,
- dataLayoutImporter.getDataLayoutSpec());
+ mlirModule->setDiscardableAttr(DLTIDialect::kDataLayoutAttrName,
+ dataLayoutImporter.getDataLayoutSpec());
return success();
}
void ModuleImport::convertTargetTriple() {
- mlirModule->setAttr(
+ mlirModule->setDiscardableAttr(
LLVM::LLVMDialect::getTargetTripleAttrName(),
builder.getStringAttr(llvmModule->getTargetTriple().str()));
}
@@ -1201,8 +1204,8 @@ void ModuleImport::convertModuleLevelAsm() {
asmArrayAttr.push_back(builder.getStringAttr(line));
}
- mlirModule->setAttr(LLVM::LLVMDialect::getModuleLevelAsmAttrName(),
- builder.getArrayAttr(asmArrayAttr));
+ mlirModule->setDiscardableAttr(LLVM::LLVMDialect::getModuleLevelAsmAttrName(),
+ builder.getArrayAttr(asmArrayAttr));
}
LogicalResult ModuleImport::convertFunctions() {
@@ -1284,7 +1287,7 @@ void ModuleImport::setFastmathFlagsAttr(llvm::Instruction *inst,
value = bitEnumSet(value, FastmathFlags::afn, flags.approxFunc());
value = bitEnumSet(value, FastmathFlags::reassoc, flags.allowReassoc());
FastmathFlagsAttr attr = FastmathFlagsAttr::get(builder.getContext(), value);
- iface->setAttr(iface.getFastmathAttrName(), attr);
+ iface.setFastmathAttr(attr);
}
/// Returns `type` if it is a builtin integer or floating-point vector type that
diff --git a/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp b/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
index 5bb42575a0955..fb70107989adc 100644
--- a/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
+++ b/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
@@ -949,9 +949,11 @@ llvm::CallInst *mlir::LLVM::detail::createIntrinsicCall(
SmallVector<llvm::OperandBundleDef> opBundles;
size_t numOpBundleOperands = 0;
auto opBundleSizesAttr = cast_if_present<DenseI32ArrayAttr>(
- intrOp->getAttr(LLVMDialect::getOpBundleSizesAttrName()));
+ intrOp->getInherentAttr(LLVMDialect::getOpBundleSizesAttrName())
+ .value_or(Attribute{}));
auto opBundleTagsAttr = cast_if_present<ArrayAttr>(
- intrOp->getAttr(LLVMDialect::getOpBundleTagsAttrName()));
+ intrOp->getInherentAttr(LLVMDialect::getOpBundleTagsAttrName())
+ .value_or(Attribute{}));
if (opBundleSizesAttr && opBundleTagsAttr) {
ArrayRef<int> opBundleSizes = opBundleSizesAttr.asArrayRef();
@@ -982,7 +984,7 @@ llvm::CallInst *mlir::LLVM::detail::createIntrinsicCall(
SmallVector<llvm::Value *> args(immArgPositions.size() + operands.size());
for (auto [immArgPos, immArgName] :
llvm::zip(immArgPositions, immArgAttrNames)) {
- Attribute attr = intrOp->getAttr(immArgName);
+ Attribute attr = intrOp->getInherentAttr(immArgName).value_or(Attribute{});
if (auto intrinsicIntegerAttr =
dyn_cast<LLVM::IntrinsicIntegerAttrInterface>(attr))
attr = intrinsicIntegerAttr.getIntegerAttr();
@@ -2461,7 +2463,7 @@ LogicalResult ModuleTranslation::createTBAAMetadata() {
}
LogicalResult ModuleTranslation::createIdentMetadata() {
- if (auto attr = mlirModule->getAttrOfType<StringAttr>(
+ if (auto attr = mlirModule->getDiscardableAttrOfType<StringAttr>(
LLVMDialect::getIdentAttrName())) {
StringRef ident = attr;
llvm::LLVMContext &ctx = llvmModule->getContext();
@@ -2475,7 +2477,7 @@ LogicalResult ModuleTranslation::createIdentMetadata() {
}
LogicalResult ModuleTranslation::createCommandlineMetadata() {
- if (auto attr = mlirModule->getAttrOfType<StringAttr>(
+ if (auto attr = mlirModule->getDiscardableAttrOfType<StringAttr>(
LLVMDialect::getCommandlineAttrName())) {
StringRef cmdLine = attr;
llvm::LLVMContext &ctx = llvmModule->getContext();
diff --git a/mlir/lib/Target/LLVMIR/Transforms/TargetToDataLayout.cpp b/mlir/lib/Target/LLVMIR/Transforms/TargetToDataLayout.cpp
index c0f9ceb6313d0..1246c9cacf98f 100644
--- a/mlir/lib/Target/LLVMIR/Transforms/TargetToDataLayout.cpp
+++ b/mlir/lib/Target/LLVMIR/Transforms/TargetToDataLayout.cpp
@@ -33,7 +33,7 @@ struct TargetToDataLayoutPass
if (initializeLLVMTargets)
LLVM::detail::initializeBackendsOnce();
- auto targetAttr = op->getAttrOfType<LLVM::TargetAttrInterface>(
+ auto targetAttr = op->getDiscardableAttrOfType<LLVM::TargetAttrInterface>(
LLVM::LLVMDialect::getTargetAttrName());
if (!targetAttr) {
op->emitError()
@@ -52,11 +52,12 @@ struct TargetToDataLayoutPass
DataLayoutSpecInterface dataLayoutSpec =
mlir::translateDataLayout(dataLayout.value(), &getContext());
- if (auto existingDlSpec = op->getAttrOfType<DataLayoutSpecInterface>(
- DLTIDialect::kDataLayoutAttrName)) {
+ if (auto existingDlSpec =
+ op->getDiscardableAttrOfType<DataLayoutSpecInterface>(
+ DLTIDialect::kDataLayoutAttrName)) {
dataLayoutSpec = existingDlSpec.combineWith({dataLayoutSpec});
}
- op->setAttr(DLTIDialect::kDataLayoutAttrName, dataLayoutSpec);
+ op->setDiscardableAttr(DLTIDialect::kDataLayoutAttrName, dataLayoutSpec);
}
};
diff --git a/mlir/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp b/mlir/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp
index 24f37cce513af..0cc5f5aabb630 100644
--- a/mlir/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp
+++ b/mlir/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp
@@ -37,7 +37,7 @@ struct TargetToTargetFeaturesPass
if (initializeLLVMTargets)
LLVM::detail::initializeBackendsOnce();
- auto targetAttr = op->getAttrOfType<LLVM::TargetAttr>(
+ auto targetAttr = op->getDiscardableAttrOfType<LLVM::TargetAttr>(
LLVM::LLVMDialect::getTargetAttrName());
if (!targetAttr) {
op->emitError() << "no LLVM::TargetAttr attribute at key \""
@@ -74,6 +74,7 @@ struct TargetToTargetFeaturesPass
LLVM::TargetAttr::get(&getContext(), targetAttr.getTriple(),
targetAttr.getChip(), fullTargetFeaturesAttr);
- op->setAttr(LLVM::LLVMDialect::getTargetAttrName(), updatedTargetAttr);
+ op->setDiscardableAttr(LLVM::LLVMDialect::getTargetAttrName(),
+ updatedTargetAttr);
}
};
diff --git a/mlir/python/mlir/dialects/gpu/__init__.py b/mlir/python/mlir/dialects/gpu/__init__.py
index b75bd525c9487..2ce4e3a24bf6b 100644
--- a/mlir/python/mlir/dialects/gpu/__init__.py
+++ b/mlir/python/mlir/dialects/gpu/__init__.py
@@ -86,7 +86,10 @@ def __init__(
if not isinstance(function_type, TypeAttr)
else function_type
)
+ if not isinstance(sym_name, (str, StringAttr)):
+ raise ValueError("sym_name must be a string or a StringAttr")
super().__init__(
+ sym_name,
function_type,
arg_attrs=arg_attrs,
res_attrs=res_attrs,
@@ -96,13 +99,6 @@ def __init__(
ip=ip,
)
- if isinstance(sym_name, str):
- self.attributes[self.SYM_NAME_ATTR_NAME] = StringAttr.get(sym_name)
- elif isinstance(sym_name, StringAttr):
- self.attributes[self.SYM_NAME_ATTR_NAME] = sym_name
- else:
- raise ValueError("sym_name must be a string or a StringAttr")
-
if kernel:
self.attributes[self.KERNEL_ATTR_NAME] = UnitAttr.get()
diff --git a/mlir/test/Conversion/ArithToLLVM/attribute-storage.mlir b/mlir/test/Conversion/ArithToLLVM/attribute-storage.mlir
new file mode 100644
index 0000000000000..662128c2af1e4
--- /dev/null
+++ b/mlir/test/Conversion/ArithToLLVM/attribute-storage.mlir
@@ -0,0 +1,32 @@
+// RUN: mlir-opt -pass-pipeline="builtin.module(func.func(convert-arith-to-llvm))" \
+// RUN: -mlir-print-op-generic %s | FileCheck %s
+
+// Check that inherent attributes are converted to target properties while
+// discardable attributes remain in the attribute dictionary.
+
+func.func @attribute_storage(%f0: f32, %f1: f32, %i0: i32, %i1: i32) {
+ // CHECK: "llvm.fadd"
+ // CHECK-SAME: <{fastmathFlags = #llvm.fastmath<fast>}>
+ // CHECK-SAME: {test.discardable = 0 : i64}
+ %0 = arith.addf %f0, %f1 fastmath<fast>
+ {"test.discardable" = 0 : i64} : f32
+
+ // CHECK: "llvm.add"
+ // CHECK-SAME: <{overflowFlags = 1 : i32}>
+ // CHECK-SAME: {test.discardable = 1 : i64}
+ %1 = arith.addi %i0, %i1 overflow<nsw>
+ {"test.discardable" = 1 : i64} : i32
+
+ // CHECK: "llvm.zext"
+ // CHECK-SAME: <{nonNeg}>
+ // CHECK-SAME: {test.discardable = 2 : i64}
+ %2 = arith.extui %i0 nneg {"test.discardable" = 2 : i64} : i32 to i64
+
+ // CHECK: "llvm.intr.experimental.constrained.fadd"
+ // CHECK-SAME: <{fastmathFlags = #llvm.fastmath<none>,
+ // CHECK-SAME: fpExceptionBehavior = 0 : i64, roundingmode = 1 : i64}>
+ // CHECK-SAME: {test.discardable = 3 : i64}
+ %3 = arith.addf %f0, %f1 to_nearest_even fastmath<fast>
+ {"test.discardable" = 3 : i64} : f32
+ return
+}
diff --git a/mlir/test/lib/Dialect/LLVM/TestPatterns.cpp b/mlir/test/lib/Dialect/LLVM/TestPatterns.cpp
index 69a3d98bc09e4..2252db9bc9802 100644
--- a/mlir/test/lib/Dialect/LLVM/TestPatterns.cpp
+++ b/mlir/test/lib/Dialect/LLVM/TestPatterns.cpp
@@ -79,8 +79,9 @@ struct TestLLVMLegalizePatternsPass
ConversionTarget target(*ctx);
target.addLegalOp(OperationName("test.legal_op", ctx));
target.addLegalDialect<LLVM::LLVMDialect>();
- target.addDynamicallyLegalOp<func::FuncOp>(
- [&](func::FuncOp funcOp) { return funcOp->hasAttr("is_legal"); });
+ target.addDynamicallyLegalOp<func::FuncOp>([&](func::FuncOp funcOp) {
+ return funcOp->hasDiscardableAttr("is_legal");
+ });
// Handle a partial conversion.
DenseSet<Operation *> unlegalizedOps;
diff --git a/mlir/test/lib/Dialect/XeGPU/TestXeGPUTransforms.cpp b/mlir/test/lib/Dialect/XeGPU/TestXeGPUTransforms.cpp
index f6b0c50da91dd..a4813528e98dd 100644
--- a/mlir/test/lib/Dialect/XeGPU/TestXeGPUTransforms.cpp
+++ b/mlir/test/lib/Dialect/XeGPU/TestXeGPUTransforms.cpp
@@ -166,7 +166,7 @@ class TestStepOpPattern : public OpConversionPattern<vector::StepOp> {
ConversionPatternRewriter &rewriter) const override {
auto layoutName = xegpu::getTemporaryLayoutName(op->getResult(0));
- auto sliceAttr = op->getAttrOfType<xegpu::SliceAttr>(layoutName);
+ auto sliceAttr = op->getDiscardableAttrOfType<xegpu::SliceAttr>(layoutName);
if (!sliceAttr || sliceAttr.getRank() != 1)
return failure();
@@ -437,7 +437,8 @@ struct TestXeGPULayoutInterface
target.addDynamicallyLegalOp<vector::StepOp>(
[&](vector::StepOp op) -> bool {
auto layoutName = xegpu::getTemporaryLayoutName(op->getResult(0));
- auto sliceAttr = op->getAttrOfType<xegpu::SliceAttr>(layoutName);
+ auto sliceAttr =
+ op->getDiscardableAttrOfType<xegpu::SliceAttr>(layoutName);
return isLegal(sliceAttr);
});
diff --git a/mlir/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp b/mlir/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp
index 5271923d923e2..c2ca29932014a 100644
--- a/mlir/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp
+++ b/mlir/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp
@@ -105,7 +105,8 @@ std::optional<mlir::gpu::SerializedObject>
TargetAttrImpl::serializeToObject(Attribute attribute, Operation *module,
const gpu::TargetOptions &options) const {
// Set a dummy attr to be retrieved by `createObject`.
- module->setAttr("serialize_attr", UnitAttr::get(module->getContext()));
+ module->setDiscardableAttr("serialize_attr",
+ UnitAttr::get(module->getContext()));
std::string targetTriple = llvm::sys::getProcessTriple();
LLVM::ModuleToObject serializer(
*module, targetTriple, "", "", 3, options.getInitialLlvmIRCallback(),
@@ -124,7 +125,7 @@ TargetAttrImpl::createObject(Attribute attribute, Operation *module,
StringAttr::get(
module->getContext(),
StringRef(object.getObject().data(), object.getObject().size())),
- module->getAttrDictionary(), /*kernels=*/nullptr);
+ module->getDiscardableAttrDictionary(), /*kernels=*/nullptr);
}
// This test checks the correct functioning of `TargetAttrInterface` as an API.
More information about the Mlir-commits
mailing list