[Mlir-commits] [mlir] [mlir][LLVM][GPU] Migrate to explicit split inherent/discardable attribute APIs access (PR #218921)
Mehdi Amini
llvmlistbot at llvm.org
Wed Sep 2 13:56:49 PDT 2026
https://github.com/joker-eph updated https://github.com/llvm/llvm-project/pull/218921
>From 9c456ccca30a7bd8301f0f5b5f90ce5f6d5a4759 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 1/3] [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
(cherry picked from commit a99328392de1f34e165d624f18d78558ec485e88)
---
.../ArithCommon/AttrToLLVMConverter.h | 53 ++++++-------------
.../mlir/Conversion/LLVMCommon/Pattern.h | 5 +-
.../Conversion/LLVMCommon/VectorPattern.h | 7 ++-
mlir/include/mlir/Dialect/GPU/IR/GPUOps.td | 2 +-
.../LLVMIR/BasicPtxBuilderInterface.td | 17 +++---
.../mlir/Dialect/LLVMIR/LLVMInterfaces.td | 11 ++++
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 3 +-
mlir/include/mlir/IR/SymbolInterfaces.td | 9 +++-
.../AMDGPUToROCDL/AMDGPUToROCDL.cpp | 6 +--
.../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 | 29 +++++-----
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 +-
.../ArithToLLVM/attribute-storage.mlir | 32 +++++++++++
mlir/test/lib/Dialect/LLVM/TestPatterns.cpp | 5 +-
.../lib/Dialect/XeGPU/TestXeGPUTransforms.cpp | 5 +-
.../Target/LLVM/SerializeToLLVMBitcode.cpp | 5 +-
54 files changed, 361 insertions(+), 239 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 a5ed42fe40d23..d9464d4dce066 100644
--- a/mlir/include/mlir/Conversion/LLVMCommon/Pattern.h
+++ b/mlir/include/mlir/Conversion/LLVMCommon/Pattern.h
@@ -344,8 +344,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 13fe6d3f25c12..e6573c94395d9 100644
--- a/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -450,7 +450,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 1a6190c41a075..e7b2ad6d05c21 100644
--- a/mlir/include/mlir/IR/SymbolInterfaces.td
+++ b/mlir/include/mlir/IR/SymbolInterfaces.td
@@ -34,7 +34,10 @@ def Symbol : OpInterface<"SymbolOpInterface"> {
"::mlir::StringAttr", "getNameAttr", (ins)
>,
InterfaceMethod<"Sets the name of this symbol.",
- "void", "setName", (ins "::mlir::StringAttr":$name)
+ "void", "setName", (ins "::mlir::StringAttr":$name), [{}],
+ /*defaultImplementation=*/[{
+ mlir::SymbolTable::setSymbolName(this->getOperation(), name);
+ }]
>,
InterfaceMethod<"Gets the visibility of this symbol.",
"mlir::SymbolTable::Visibility", "getVisibility", (ins)
@@ -183,7 +186,9 @@ def Symbol : OpInterface<"SymbolOpInterface"> {
// Add additional classof checks to properly handle "optional" symbols.
let extraClassOf = [{
- return static_cast<bool>($_op.getNameAttr());
+ 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 2a846339b562e..2f6ede5d88d16 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/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 e2d251bb109f6..66322f9ebc2f6 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 9c116ef4f7ac2..bc43fe6b33dff 100644
--- a/mlir/lib/Conversion/LLVMCommon/Pattern.cpp
+++ b/mlir/lib/Conversion/LLVMCommon/Pattern.cpp
@@ -372,7 +372,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 7207c960985f3..4bc0d3a7a6ebe 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 d6ff8504582c8..662f83ce2702e 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 f3855229e8fbd..f4145967e94dc 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 49b697f906318..74bb7eb4b1665 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();
}
@@ -4203,7 +4208,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 02cdb6088b6ac..7c1b8b263b591 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();
}
@@ -1180,13 +1183,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()));
}
@@ -1202,8 +1205,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() {
@@ -1285,7 +1288,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/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.
>From 3a845492c4df1fde9d8126abb6b351fe202420ba Mon Sep 17 00:00:00 2001
From: Mehdi Amini <joker.eph at gmail.com>
Date: Thu, 27 Aug 2026 13:55:22 -0700
Subject: [PATCH 2/3] Add SymbolTable attribute-name compatibility helper
Provide the zero-argument getSymbolAttrName() accessor expected by
generated Flang and MLIR code that still uses the conventional sym_name
attribute name.
Assisted-by: Codex
(cherry picked from commit 909be360ce1192bf59924f6633cc762bb06f5f3d)
---
mlir/include/mlir/IR/SymbolTable.h | 6 ++++++
1 file changed, 6 insertions(+)
diff --git a/mlir/include/mlir/IR/SymbolTable.h b/mlir/include/mlir/IR/SymbolTable.h
index 96be32ee99467..f8af0e547fedd 100644
--- a/mlir/include/mlir/IR/SymbolTable.h
+++ b/mlir/include/mlir/IR/SymbolTable.h
@@ -118,6 +118,12 @@ class SymbolTable {
/// Returns the name of the given symbol operation, aborting if no symbol is
/// present.
static StringAttr getSymbolName(Operation *symbol);
+ /// Compatibility alias for generated interface code that still refers to the
+ /// symbol name attribute by convention.
+ static StringRef getSymbolAttrName() { return "sym_name"; }
+ static StringAttr getSymbolAttrName(Operation *symbol) {
+ return getSymbolName(symbol);
+ }
/// Sets the name of the given symbol operation.
static void setSymbolName(Operation *symbol, StringAttr name);
>From 84d2ffc4829e3b755c3421658cd947b0e59f456e Mon Sep 17 00:00:00 2001
From: Mehdi Amini <joker.eph at gmail.com>
Date: Thu, 27 Aug 2026 14:40:39 -0700
Subject: [PATCH 3/3] [MLIR][LLVM][GPU] Fix explicit attribute migration issues
Fix the SymbolOpInterface setter collision without recursive dispatch and
retain deprecated forwarding setters for source compatibility.
Preserve inherent and discardable attribute storage across conversion,
printing, target feature lookup, and XeGPU rewrites. Convert floating-point
extension and truncation fast-math properties across dialect schemas, and drop
stale XeGPU instruction data from typed layout properties.
Add regression coverage for ArithToLLVM, MathToXeVM, LLVM target features,
symbol helpers, and XeGPU blocking.
Assisted-by: Codex
---
.../XeGPU/Transforms/XeGPULayoutImpl.h | 4 ++
mlir/include/mlir/IR/SymbolInterfaces.td | 24 +++++---
mlir/include/mlir/IR/SymbolTable.h | 3 -
.../Conversion/ArithToLLVM/ArithToLLVM.cpp | 9 +--
.../ControlFlowToLLVM/ControlFlowToLLVM.cpp | 4 +-
mlir/lib/Conversion/LLVMCommon/Pattern.cpp | 4 +-
mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp | 13 ++++-
mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp | 4 +-
mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp | 18 +++++-
mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp | 2 +-
.../XeGPU/Transforms/XeGPUBlocking.cpp | 1 +
.../XeGPU/Transforms/XeGPULayoutImpl.cpp | 7 +++
.../XeGPU/Transforms/XeGPUPropagateLayout.cpp | 9 ++-
.../Transforms/XeGPUSgToLaneDistribute.cpp | 11 ++--
.../Dialect/XeGPU/Transforms/XeGPUUnroll.cpp | 15 +++--
mlir/lib/IR/SymbolTable.cpp | 2 +-
.../ArithToLLVM/attribute-storage.mlir | 8 +++
.../Conversion/MathToXeVM/math-to-xevm.mlir | 3 +
mlir/test/Dialect/XeGPU/xegpu-blocking.mlir | 12 ++--
mlir/test/lib/Dialect/Test/TestOpDefs.cpp | 4 --
mlir/unittests/Dialect/LLVMIR/CMakeLists.txt | 2 +
.../Dialect/LLVMIR/LLVMAttrsTest.cpp | 55 +++++++++++++++++++
mlir/unittests/IR/SymbolTableTest.cpp | 2 +-
23 files changed, 162 insertions(+), 54 deletions(-)
create mode 100644 mlir/unittests/Dialect/LLVMIR/LLVMAttrsTest.cpp
diff --git a/mlir/include/mlir/Dialect/XeGPU/Transforms/XeGPULayoutImpl.h b/mlir/include/mlir/Dialect/XeGPU/Transforms/XeGPULayoutImpl.h
index defc389006af5..9d0ee049ad936 100644
--- a/mlir/include/mlir/Dialect/XeGPU/Transforms/XeGPULayoutImpl.h
+++ b/mlir/include/mlir/Dialect/XeGPU/Transforms/XeGPULayoutImpl.h
@@ -93,6 +93,10 @@ dropSgLayoutAndDataOnAttrs(ArrayRef<NamedAttribute> attrs);
/// any DistributeLayoutAttr found.
SmallVector<NamedAttribute> dropInstDataOnAttrs(ArrayRef<NamedAttribute> attrs);
+/// Drops inst-data information from DistributeLayoutAttrs stored as inherent
+/// attributes on the operation.
+void dropInstDataOnInherentAttrs(Operation *op);
+
//===----------------------------------------------------------------------===//
// Backward layout inference (result layout -> source layout)
//===----------------------------------------------------------------------===//
diff --git a/mlir/include/mlir/IR/SymbolInterfaces.td b/mlir/include/mlir/IR/SymbolInterfaces.td
index e7b2ad6d05c21..97570b19de6f9 100644
--- a/mlir/include/mlir/IR/SymbolInterfaces.td
+++ b/mlir/include/mlir/IR/SymbolInterfaces.td
@@ -34,9 +34,11 @@ def Symbol : OpInterface<"SymbolOpInterface"> {
"::mlir::StringAttr", "getNameAttr", (ins)
>,
InterfaceMethod<"Sets the name of this symbol.",
- "void", "setName", (ins "::mlir::StringAttr":$name), [{}],
+ "void", "setSymbolName", (ins "::mlir::StringAttr":$name), [{}],
/*defaultImplementation=*/[{
- mlir::SymbolTable::setSymbolName(this->getOperation(), name);
+ auto *op = this->getOperation();
+ op->setInherentAttr(::mlir::StringAttr::get(
+ op->getContext(), mlir::SymbolTable::getSymbolAttrName()), name);
}]
>,
InterfaceMethod<"Gets the visibility of this symbol.",
@@ -138,9 +140,19 @@ def Symbol : OpInterface<"SymbolOpInterface"> {
return getNameAttr().getValue();
}
- /// Convenience version of `setName` that takes a StringRef.
+ /// Convenience version of `setSymbolName` that takes a StringRef.
+ void setSymbolName(::mlir::StringRef name) {
+ setSymbolName(::mlir::StringAttr::get(getOperation()->getContext(), name));
+ }
+
+ [[deprecated("use setSymbolName instead")]]
+ void setName(::mlir::StringAttr name) {
+ setSymbolName(name);
+ }
+
+ [[deprecated("use setSymbolName instead")]]
void setName(::mlir::StringRef name) {
- setName(::mlir::StringAttr::get(getOperation()->getContext(), name));
+ setSymbolName(name);
}
/// Return the conventional symbol visibility attribute name.
@@ -186,9 +198,7 @@ def Symbol : OpInterface<"SymbolOpInterface"> {
// Add additional classof checks to properly handle "optional" symbols.
let extraClassOf = [{
- return static_cast<bool>(
- $_op->getInherentAttr(::mlir::SymbolTable::getSymbolAttrName())
- .value_or(::mlir::Attribute{}));
+ return static_cast<bool>($_op.getNameAttr());
}];
}
diff --git a/mlir/include/mlir/IR/SymbolTable.h b/mlir/include/mlir/IR/SymbolTable.h
index f8af0e547fedd..a6d8e8660cd22 100644
--- a/mlir/include/mlir/IR/SymbolTable.h
+++ b/mlir/include/mlir/IR/SymbolTable.h
@@ -121,9 +121,6 @@ class SymbolTable {
/// Compatibility alias for generated interface code that still refers to the
/// symbol name attribute by convention.
static StringRef getSymbolAttrName() { return "sym_name"; }
- static StringAttr getSymbolAttrName(Operation *symbol) {
- return getSymbolName(symbol);
- }
/// Sets the name of the given symbol operation.
static void setSymbolName(Operation *symbol, StringAttr name);
diff --git a/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp b/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
index 85cf4275e41a8..c13d7ae16e0ad 100644
--- a/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
+++ b/mlir/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
@@ -113,9 +113,10 @@ using DivSIOpLowering =
VectorConvertToLLVMPattern<arith::DivSIOp, LLVM::SDivOp>;
using DivUIOpLowering =
VectorConvertToLLVMPattern<arith::DivUIOp, LLVM::UDivOp>;
-using ExtFOpLowering = VectorConvertToLLVMPattern<arith::ExtFOp, LLVM::FPExtOp,
- AttrConvertPassThrough,
- /*FailOnUnsupportedFP=*/true>;
+using ExtFOpLowering =
+ VectorConvertToLLVMPattern<arith::ExtFOp, LLVM::FPExtOp,
+ arith::AttrConvertFastMathToLLVM,
+ /*FailOnUnsupportedFP=*/true>;
using ExtSIOpLowering =
VectorConvertToLLVMPattern<arith::ExtSIOp, LLVM::SExtOp>;
using ExtUIOpLowering =
@@ -202,7 +203,7 @@ using SubIOpLowering =
using TruncFOpLowering =
ConstrainedVectorConvertToLLVMPattern<arith::TruncFOp, LLVM::FPTruncOp,
/*HasRoundingMode=*/false,
- AttrConvertPassThrough,
+ arith::AttrConvertFastMathToLLVM,
/*FailOnUnsupportedFP=*/true>;
using ConstrainedTruncFOpLowering = ConstrainedVectorConvertToLLVMPattern<
arith::TruncFOp, LLVM::ConstrainedFPTruncIntr, /*HasRoundingMode=*/true,
diff --git a/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp b/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
index c09edbc7cb317..2c2b6b567bffc 100644
--- a/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
+++ b/mlir/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
@@ -135,12 +135,14 @@ static SmallVector<Value> flattenValues(ArrayRef<ValueRange> values) {
/// Set attributes on an operation using its inherent/discardable split.
static void setConvertedAttrs(Operation *op, DictionaryAttr attrs) {
+ SmallVector<NamedAttribute> discardableAttrs;
for (NamedAttribute attr : attrs) {
if (op->getInherentAttr(attr.getName()).has_value())
op->setInherentAttr(attr.getName(), attr.getValue());
else
- op->setDiscardableAttr(attr.getName(), attr.getValue());
+ discardableAttrs.push_back(attr);
}
+ op->setDiscardableAttrs(discardableAttrs);
}
/// Convert the destination block signature (if necessary) and lower the branch
diff --git a/mlir/lib/Conversion/LLVMCommon/Pattern.cpp b/mlir/lib/Conversion/LLVMCommon/Pattern.cpp
index bc43fe6b33dff..2c4b8db4adc51 100644
--- a/mlir/lib/Conversion/LLVMCommon/Pattern.cpp
+++ b/mlir/lib/Conversion/LLVMCommon/Pattern.cpp
@@ -372,17 +372,19 @@ LogicalResult LLVM::detail::intrinsicRewrite(
auto callIntrOp = LLVM::CallIntrinsicOp::create(
rewriter, loc, resType, rewriter.getStringAttr(intrinsic), operands);
// Propagate attributes.
+ SmallVector<NamedAttribute> discardableAttrs;
auto copyAttr = [&](StringAttr name, Attribute attr) {
if (callIntrOp->getInherentAttr(name).has_value())
callIntrOp->setInherentAttr(name, attr);
else
- callIntrOp->setDiscardableAttr(name, attr);
+ discardableAttrs.emplace_back(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);
});
+ callIntrOp->setDiscardableAttrs(discardableAttrs);
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 4b29df764be71..58b7948892db4 100644
--- a/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
+++ b/mlir/lib/Conversion/MathToXeVM/MathToXeVM.cpp
@@ -144,9 +144,16 @@ struct ConvertNativeFuncPattern final : public OpConversionPattern<Op> {
// calls, in order to allow further fastmath optimizations: We thus need to
// convert arith fastmath attrs into attrs recognized by llvm.
arith::AttrConvertFastMathToLLVM<Op, LLVM::CallOp> fastAttrConverter(op);
- mlir::NamedAttribute fastAttr = fastAttrConverter.getAttrs()[0];
- callOp.setFastmathFlagsAttr(
- cast<LLVM::FastmathFlagsAttr>(fastAttr.getValue()));
+ SmallVector<NamedAttribute> discardableAttrs;
+ for (mlir::NamedAttribute attr : fastAttrConverter.getAttrs()) {
+ if (attr.getName() == LLVM::CallOp::getFastmathAttrName()) {
+ callOp.setFastmathFlagsAttr(
+ cast<LLVM::FastmathFlagsAttr>(attr.getValue()));
+ continue;
+ }
+ discardableAttrs.push_back(attr);
+ }
+ callOp->setDiscardableAttrs(discardableAttrs);
if (unwrapSizeOneVec) {
// Re-wrap the scalar result back into a size-1 vector to preserve types.
diff --git a/mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp b/mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp
index 3b350ba8b6857..9af8ee71eb845 100644
--- a/mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp
+++ b/mlir/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp
@@ -551,11 +551,12 @@ static LLVM::CallOp createDeviceFunctionCall(
funcOp.setArgAttr(idx, attrName, rewriter.getUnitAttr());
auto callOp = LLVM::CallOp::create(rewriter, loc, funcOp, args);
+ SmallVector<NamedAttribute> discardableAttrs;
auto copyAttr = [&](StringAttr name, Attribute attr) {
if (callOp->getInherentAttr(name).has_value())
callOp->setInherentAttr(name, attr);
else
- callOp->setDiscardableAttr(name, attr);
+ discardableAttrs.emplace_back(name, attr);
};
for (NamedAttribute attr : funcOp->getDiscardableAttrDictionary())
copyAttr(attr.getName(), attr.getValue());
@@ -563,6 +564,7 @@ static LLVM::CallOp createDeviceFunctionCall(
funcOp, [&](StringRef name, Attribute &attr) {
copyAttr(rewriter.getStringAttr(name), attr);
});
+ callOp->setDiscardableAttrs(discardableAttrs);
return callOp;
}
diff --git a/mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp b/mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp
index fdfae67bce47a..317d6b607e08b 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp
@@ -620,8 +620,22 @@ TargetFeaturesAttr TargetFeaturesAttr::featuresAt(Operation *op) {
auto parentFunction = op->getParentOfType<FunctionOpInterface>();
if (!parentFunction)
return {};
- return parentFunction.getOperation()
- ->getDiscardableAttrOfType<TargetFeaturesAttr>(getAttributeName());
+ Operation *parentOp = parentFunction.getOperation();
+ if (auto llvmFunc = dyn_cast<LLVMFuncOp>(parentOp))
+ return llvmFunc.getTargetFeaturesAttr();
+
+ // Dynamic operation definitions do not support direct inherent attribute
+ // lookup. Use the safe walker for other FunctionOpInterface implementations.
+ TargetFeaturesAttr targetFeatures;
+ parentOp->getName().walkInherentAttrs(
+ parentOp, [&](StringRef name, Attribute &attr) {
+ if (name == getAttributeName())
+ targetFeatures = dyn_cast<TargetFeaturesAttr>(attr);
+ });
+ if (targetFeatures)
+ return targetFeatures;
+ return parentOp->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 74bb7eb4b1665..47cc0f64fe906 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
@@ -54,7 +54,7 @@ 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()));
+ return attrs;
}
static auto processFMFAttr(ArrayRef<NamedAttribute> attrs) {
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp
index f70ff26b3d8f7..d9648a6e04d1f 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp
@@ -545,6 +545,7 @@ void XeGPUBlockingPass::runOnOperation() {
SmallVector<NamedAttribute> newAttrs = xegpu::dropInstDataOnAttrs(
op->getDiscardableAttrDictionary().getValue());
op->setDiscardableAttrs(newAttrs);
+ xegpu::dropInstDataOnInherentAttrs(op);
});
// 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 4825d12d347c4..24374239461d1 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp
@@ -70,6 +70,13 @@ xegpu::dropInstDataOnAttrs(ArrayRef<NamedAttribute> attrs) {
return out;
}
+void xegpu::dropInstDataOnInherentAttrs(Operation *op) {
+ op->getName().walkInherentAttrs(op, [](StringRef, Attribute &attr) {
+ if (auto dist = dyn_cast<xegpu::DistributeLayoutAttr>(attr))
+ attr = dist.dropInstData();
+ });
+}
+
// Sets the layout on a TensorDesc value by updating its type to include
// the given layout, if the type does not already have a layout attached.
static void setTensorDescLayout(Value val, xegpu::DistributeLayoutAttr layout) {
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp
index b7f44f07804c3..a94b06ec4ee05 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp
@@ -1753,12 +1753,11 @@ ResolveLayoutConflicts::resolveTensorDescConsumer(OpOperand &operand) {
conflictingCreateNdOp.getContext(), currTDescType.getShape(),
currTDescType.getElementType(), currTDescType.getEncoding(),
expectedLayout);
- OperationState state(
- consumerOp->getLoc(), conflictingCreateNdOp->getName(),
- conflictingCreateNdOp->getOperands(), TypeRange{newTensorDescType},
+ auto newOp = xegpu::CreateNdDescOp::create(
+ builder, consumerOp->getLoc(), TypeRange{newTensorDescType},
+ conflictingCreateNdOp->getOperands(),
+ conflictingCreateNdOp.getProperties(),
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 b84fe9aa97b5c..812233047fb88 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp
@@ -127,12 +127,11 @@ struct SgToLaneCreateNdDesc
if (!resultType.getLayout())
return failure();
- 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));
+ auto newOp = xegpu::CreateNdDescOp::create(
+ rewriter, op.getLoc(), TypeRange{resultType.dropLayouts()},
+ op.getOperands(), op.getProperties(),
+ op->getDiscardableAttrDictionary().getValue());
+ rewriter.replaceOp(op, newOp.getResult());
return success();
}
};
diff --git a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp
index 2be5d21e5f46d..1854c639c4797 100644
--- a/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp
+++ b/mlir/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp
@@ -428,12 +428,12 @@ struct UnrollDpasOp : public UnrollPattern<xegpu::DpasOp> {
if (tmpC)
operands.push_back(tmpC);
- OperationState state(
- loc, op->getName(), operands, TypeRange{vecTy},
+ auto newDpasOp = xegpu::DpasOp::create(
+ rewriter, loc, TypeRange{vecTy}, operands, op.getProperties(),
xegpu::dropInstDataOnAttrs(
op->getDiscardableAttrDictionary().getValue()));
- state.propertiesAttr = op->getPropertiesAsAttribute();
- tmpC = rewriter.create(state)->getResult(0);
+ xegpu::dropInstDataOnInherentAttrs(newDpasOp);
+ tmpC = newDpasOp.getResult();
}
newOps.push_back(tmpC);
}
@@ -533,12 +533,11 @@ struct UnrollDpasMxOp : public UnrollPattern<xegpu::DpasMxOp> {
operands.push_back(
bScaleVals[batch * (kIters * nIters) + k * nIters + j]);
- OperationState state(
- loc, op->getName(), operands, TypeRange{vecTy},
+ newDpasMxOp = xegpu::DpasMxOp::create(
+ rewriter, loc, TypeRange{vecTy}, operands, op.getProperties(),
xegpu::dropInstDataOnAttrs(
op->getDiscardableAttrDictionary().getValue()));
- state.propertiesAttr = op->getPropertiesAsAttribute();
- newDpasMxOp = cast<xegpu::DpasMxOp>(rewriter.create(state));
+ xegpu::dropInstDataOnInherentAttrs(newDpasMxOp);
tmpC = newDpasMxOp.getResult();
}
newOps.push_back(newDpasMxOp);
diff --git a/mlir/lib/IR/SymbolTable.cpp b/mlir/lib/IR/SymbolTable.cpp
index 31e9ae5c5dfe2..b0288185ae0ab 100644
--- a/mlir/lib/IR/SymbolTable.cpp
+++ b/mlir/lib/IR/SymbolTable.cpp
@@ -295,7 +295,7 @@ StringAttr SymbolTable::getSymbolName(Operation *symbol) {
/// Sets the name of the given symbol operation.
void SymbolTable::setSymbolName(Operation *symbol, StringAttr name) {
auto symbolOp = cast<SymbolOpInterface>(symbol);
- symbolOp.setName(name);
+ symbolOp.setSymbolName(name);
}
/// Returns the visibility of the given symbol operation.
diff --git a/mlir/test/Conversion/ArithToLLVM/attribute-storage.mlir b/mlir/test/Conversion/ArithToLLVM/attribute-storage.mlir
index 662128c2af1e4..5ccc0f0e609b8 100644
--- a/mlir/test/Conversion/ArithToLLVM/attribute-storage.mlir
+++ b/mlir/test/Conversion/ArithToLLVM/attribute-storage.mlir
@@ -28,5 +28,13 @@ func.func @attribute_storage(%f0: f32, %f1: f32, %i0: i32, %i1: i32) {
// CHECK-SAME: {test.discardable = 3 : i64}
%3 = arith.addf %f0, %f1 to_nearest_even fastmath<fast>
{"test.discardable" = 3 : i64} : f32
+
+ // CHECK: "llvm.fpext"
+ // CHECK-SAME: <{fastmathFlags = #llvm.fastmath<nnan>}>
+ %4 = arith.extf %f0 fastmath<nnan> : f32 to f64
+
+ // CHECK: "llvm.fptrunc"
+ // CHECK-SAME: <{fastmathFlags = #llvm.fastmath<nnan>}>
+ %5 = arith.truncf %f0 fastmath<nnan> : f32 to f16
return
}
diff --git a/mlir/test/Conversion/MathToXeVM/math-to-xevm.mlir b/mlir/test/Conversion/MathToXeVM/math-to-xevm.mlir
index 0ca8d14096d05..55eb300b70994 100644
--- a/mlir/test/Conversion/MathToXeVM/math-to-xevm.mlir
+++ b/mlir/test/Conversion/MathToXeVM/math-to-xevm.mlir
@@ -67,6 +67,9 @@ module @test_module {
// CHECK: llvm.call @_Z22__spirv_ocl_native_expd(%{{.*}}) {fastmathFlags = #llvm.fastmath<afn>} : (f64) -> f64
%exp_afn_f64 = math.exp %c1_f64 fastmath<afn> : f64
+ // CHECK: llvm.call @_Z22__spirv_ocl_native_expd(%{{.*}}) {aaa = 0 : i32, fastmathFlags = #llvm.fastmath<afn>} : (f64) -> f64
+ %exp_with_discardable_attr = math.exp %c1_f64 fastmath<afn> {aaa = 0 : i32} : f64
+
// CHECK: math.exp
%exp_none_f16 = math.exp %c1_f16 fastmath<none> : f16
// CHECK: math.exp
diff --git a/mlir/test/Dialect/XeGPU/xegpu-blocking.mlir b/mlir/test/Dialect/XeGPU/xegpu-blocking.mlir
index 44b7615ac9af7..c7cbdf16a1171 100644
--- a/mlir/test/Dialect/XeGPU/xegpu-blocking.mlir
+++ b/mlir/test/Dialect/XeGPU/xegpu-blocking.mlir
@@ -26,7 +26,7 @@ gpu.module @test_kernel {
%a = xegpu.load_nd %a_tdesc[%c0, %k] <{layout = #a}>: !xegpu.tensor_desc<16x32xf16, #a> -> vector<16x32xf16>
//CHECK-COUNT-4: xegpu.load_nd {{.*}} -> vector<16x16xf16>
%b = xegpu.load_nd %b_tdesc[%k, %c0] <{layout = #b}>: !xegpu.tensor_desc<32x32xf16, #b> -> vector<32x32xf16>
- //CHECK-COUNT-8: xegpu.dpas {{.*}}
+ //CHECK-COUNT-8: xegpu.dpas {{.*}} <{layout_a = #xegpu.layout<lane_layout =
%c = xegpu.dpas %a, %b, %arg2 <{layout_a = #a, layout_b = #b, layout_cd = #c}>: vector<16x32xf16>, vector<32x32xf16>, vector<16x32xf32> -> vector<16x32xf32>
scf.yield %c
: vector<16x32xf32>
@@ -63,7 +63,7 @@ gpu.module @test_kernel {
%a = xegpu.load_nd %a_tdesc[%c0, %k] <{layout = #l1}>: !xegpu.tensor_desc<16x32xf16, #l1> -> vector<16x32xf16>
//CHECK-COUNT-4: xegpu.load_nd {{.*}} -> vector<16x16xf16>
%b = xegpu.load_nd %b_tdesc[%k, %c0] <{layout = #l2}>: !xegpu.tensor_desc<32x32xf16, #l2> -> vector<32x32xf16>
- //CHECK-COUNT-8: xegpu.dpas {{.*}}
+ //CHECK-COUNT-8: xegpu.dpas {{[^<]*}} :
%c = xegpu.dpas %a, %b, %arg2 <{layout_a = #l1, layout_b = #l2, layout_cd = #l1}>: vector<16x32xf16>, vector<32x32xf16>, vector<16x32xf32> -> vector<16x32xf32>
scf.yield %c
: vector<16x32xf32>
@@ -144,7 +144,7 @@ gpu.module @test_kernel {
%b = xegpu.load_nd %b_tdesc[%k, %c0] <{layout = #b}>: !xegpu.tensor_desc<32x32xf16, #b> -> vector<32x32xf16>
//CHECK-COUNT-4: math.exp {{.*}} : vector<8x16xf16>
%e = math.exp %a : vector<16x32xf16>
- //CHECK-COUNT-8: xegpu.dpas {{.*}}
+ //CHECK-COUNT-8: xegpu.dpas {{.*}} <{layout_a = #xegpu.layout<lane_layout =
%c = xegpu.dpas %e, %b, %arg2 <{layout_a = #a, layout_b = #b, layout_cd = #c}>: vector<16x32xf16>, vector<32x32xf16>, vector<16x32xf32> -> vector<16x32xf32>
scf.yield %c
: vector<16x32xf32>
@@ -780,7 +780,7 @@ gpu.module @test_kernel {
%sa = xegpu.load_nd %scale_a_tdesc[%c0, %c0] <{layout = #l1_scale}>: !xegpu.tensor_desc<16x2xf8E8M0FNU, #l1_scale> -> vector<16x2xf8E8M0FNU>
//CHECK-COUNT-4: xegpu.load_nd {{.*}} -> vector<1x16xf8E8M0FNU>
%sb = xegpu.load_nd %scale_b_tdesc[%c0, %c0] <{layout = #l2_scale}>: !xegpu.tensor_desc<2x32xf8E8M0FNU, #l2_scale> -> vector<2x32xf8E8M0FNU>
- //CHECK-COUNT-8: xegpu.dpas_mx {{.*}}
+ //CHECK-COUNT-8: xegpu.dpas_mx {{[^<]*}} :
%c = xegpu.dpas_mx %a, %b, %arg2 scale_a = %sa scale_b = %sb <{layout_a=#l1, layout_b = #l2, layout_cd = #l3, layout_a_scale = #l1_scale, layout_b_scale = #l2_scale}> {layout_result_0 = #l3}: (vector<16x64xf4E2M1FN>, vector<64x32xf4E2M1FN>, vector<16x32xf32>, vector<16x2xf8E8M0FNU>, vector<2x32xf8E8M0FNU>) -> vector<16x32xf32>
scf.yield %c : vector<16x32xf32>
} {layout_result_0 = #l3}
@@ -808,7 +808,7 @@ gpu.module @test_kernel {
%a = xegpu.load_nd %a_tdesc[0, 0, 0] <{layout = #a_3d}>: !xegpu.tensor_desc<4x8x32xf16, #a_3d> -> vector<4x8x32xf16>
// CHECK-COUNT-8: xegpu.load_nd {{.*}} -> vector<1x16x16xf16>
%b = xegpu.load_nd %b_tdesc[0, 0, 0] <{layout = #b_3d}>: !xegpu.tensor_desc<4x32x16xf16, #b_3d> -> vector<4x32x16xf16>
- // CHECK-COUNT-8: xegpu.dpas {{.*}} : vector<1x8x16xf16>, vector<1x16x16xf16>, vector<1x8x16xf32> -> vector<1x8x16xf32>
+ // CHECK-COUNT-8: xegpu.dpas {{[^<]*}} : vector<1x8x16xf16>, vector<1x16x16xf16>, vector<1x8x16xf32> -> vector<1x8x16xf32>
%d = xegpu.dpas %a, %b, %c_init <{layout_a = #a_3d, layout_b = #b_3d, layout_cd = #c_3d}>
: vector<4x8x32xf16>, vector<4x32x16xf16>, vector<4x8x16xf32> -> vector<4x8x16xf32>
// CHECK-COUNT-4: xegpu.store_nd {{.*}} : vector<1x8x16xf32>, !xegpu.tensor_desc<1x8x16xf32>
@@ -845,7 +845,7 @@ gpu.module @test_kernel {
%sb = xegpu.load_nd %sb_tdesc[0, 0, 0] <{layout = #sb_3d}>: !xegpu.tensor_desc<2x2x32xf8E8M0FNU, #sb_3d> -> vector<2x2x32xf8E8M0FNU>
// dpas_mx: [2,16,64] x [2,64,32] -> [2,16,32] with scales [2,16,2] and [2,2,32]
// unrolled: batch=2, M=16/8=2, K=64/64=1, N=32/16=2 -> 2*2*2=8 results (with k-reduction)
- // CHECK-COUNT-8: xegpu.dpas_mx {{.*}} : (vector<1x8x64xf4E2M1FN>, vector<1x64x16xf4E2M1FN>, vector<1x8x16xf32>, vector<1x8x2xf8E8M0FNU>, vector<1x2x16xf8E8M0FNU>) -> vector<1x8x16xf32>
+ // CHECK-COUNT-8: xegpu.dpas_mx {{[^<]*}} : (vector<1x8x64xf4E2M1FN>, vector<1x64x16xf4E2M1FN>, vector<1x8x16xf32>, vector<1x8x2xf8E8M0FNU>, vector<1x2x16xf8E8M0FNU>) -> vector<1x8x16xf32>
%d = xegpu.dpas_mx %a, %b, %c_init scale_a = %sa scale_b = %sb
<{layout_a = #a_3d_mx, layout_b = #b_3d_mx, layout_cd = #c_3d_mx,
layout_a_scale = #sa_3d, layout_b_scale = #sb_3d}>
diff --git a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
index 06871fd0ea835..4c683358e71a5 100644
--- a/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
+++ b/mlir/test/lib/Dialect/Test/TestOpDefs.cpp
@@ -30,10 +30,6 @@ StringAttr OverriddenSymbolVisibilityOp::getNameAttr() {
return getSymNameAttr();
}
-void OverriddenSymbolVisibilityOp::setName(StringAttr name) {
- setSymNameAttr(name);
-}
-
static StringLiteral getVisibilityString(SymbolTable::Visibility visibility) {
switch (visibility) {
case SymbolTable::Visibility::Private:
diff --git a/mlir/unittests/Dialect/LLVMIR/CMakeLists.txt b/mlir/unittests/Dialect/LLVMIR/CMakeLists.txt
index 7cc130d02ad74..dc0e82d7591e8 100644
--- a/mlir/unittests/Dialect/LLVMIR/CMakeLists.txt
+++ b/mlir/unittests/Dialect/LLVMIR/CMakeLists.txt
@@ -1,7 +1,9 @@
add_mlir_unittest(MLIRLLVMIRTests
+ LLVMAttrsTest.cpp
LLVMTypeTest.cpp
)
mlir_target_link_libraries(MLIRLLVMIRTests
PRIVATE
+ MLIRFuncDialect
MLIRLLVMDialect
)
diff --git a/mlir/unittests/Dialect/LLVMIR/LLVMAttrsTest.cpp b/mlir/unittests/Dialect/LLVMIR/LLVMAttrsTest.cpp
new file mode 100644
index 0000000000000..15b5464be7262
--- /dev/null
+++ b/mlir/unittests/Dialect/LLVMIR/LLVMAttrsTest.cpp
@@ -0,0 +1,55 @@
+//===- LLVMAttrsTest.cpp - Tests for LLVM attributes ----------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "LLVMTestBase.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/IR/Builders.h"
+#include "mlir/IR/BuiltinOps.h"
+
+using namespace mlir;
+using namespace mlir::LLVM;
+
+TEST_F(LLVMIRTest, TargetFeaturesAtLLVMFuncOp) {
+ OpBuilder builder(&context);
+ Location loc = builder.getUnknownLoc();
+ OwningOpRef<ModuleOp> module = ModuleOp::create(loc);
+ builder.setInsertionPointToStart(module->getBody());
+
+ auto functionType =
+ LLVMFunctionType::get(LLVMVoidType::get(&context), /*params=*/{});
+ LLVMFuncOp function = LLVMFuncOp::create(builder, loc, "test", functionType);
+ TargetFeaturesAttr targetFeatures = TargetFeaturesAttr::get(&context, "+sve");
+ function.setTargetFeaturesAttr(targetFeatures);
+
+ Block *entry = function.addEntryBlock(builder);
+ builder.setInsertionPointToStart(entry);
+ ReturnOp returnOp = ReturnOp::create(builder, loc, ValueRange{});
+
+ EXPECT_EQ(TargetFeaturesAttr::featuresAt(returnOp), targetFeatures);
+}
+
+TEST_F(LLVMIRTest, TargetFeaturesAtFuncOpDiscardableAttr) {
+ context.loadDialect<func::FuncDialect>();
+ OpBuilder builder(&context);
+ Location loc = builder.getUnknownLoc();
+ OwningOpRef<ModuleOp> module = ModuleOp::create(loc);
+ builder.setInsertionPointToStart(module->getBody());
+
+ auto functionType = builder.getFunctionType(/*inputs=*/{}, /*results=*/{});
+ func::FuncOp function =
+ func::FuncOp::create(builder, loc, "test", functionType);
+ TargetFeaturesAttr targetFeatures = TargetFeaturesAttr::get(&context, "+sve");
+ function->setDiscardableAttr(TargetFeaturesAttr::getAttributeName(),
+ targetFeatures);
+
+ Block *entry = function.addEntryBlock();
+ builder.setInsertionPointToStart(entry);
+ func::ReturnOp returnOp = func::ReturnOp::create(builder, loc);
+
+ EXPECT_EQ(TargetFeaturesAttr::featuresAt(returnOp), targetFeatures);
+}
diff --git a/mlir/unittests/IR/SymbolTableTest.cpp b/mlir/unittests/IR/SymbolTableTest.cpp
index f18cf85ef0568..fd96467e7c76e 100644
--- a/mlir/unittests/IR/SymbolTableTest.cpp
+++ b/mlir/unittests/IR/SymbolTableTest.cpp
@@ -89,7 +89,7 @@ TEST(SymbolOpInterface, NativeSymbolTraits) {
EXPECT_EQ(symOp.getName(), "symbol_name");
EXPECT_TRUE(symOp.isPublic());
- symOp.setName("new_name");
+ symOp.setSymbolName("new_name");
EXPECT_EQ(symOp.getName(), "new_name");
EXPECT_EQ(symOp->getInherentAttr("sym_name").value_or(Attribute{}),
symOp.getNameAttr());
More information about the Mlir-commits
mailing list