[mlir][LLVM][GPU] Migrate to explicit split inherent/discardable attribute APIs access (#218921) Use discardable attribute APIs and typed operation accessors throughout the LLVM and GPU dialect families, their conversions, translations, and tests. Assisted-by: Codex GitOrigin-RevId: 96cb9c493b96f1ddf0c47c3bb7c1195b88148ec7
diff --git a/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h b/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h index feb74c8..cdea530 100644 --- a/include/mlir/Conversion/ArithCommon/AttrToLLVMConverter.h +++ b/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 @@ 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 @@ 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 @@ 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 @@ "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/include/mlir/Conversion/LLVMCommon/Pattern.h b/include/mlir/Conversion/LLVMCommon/Pattern.h index a5ed42f..d9464d4 100644 --- a/include/mlir/Conversion/LLVMCommon/Pattern.h +++ b/include/mlir/Conversion/LLVMCommon/Pattern.h
@@ -344,8 +344,9 @@ 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/include/mlir/Conversion/LLVMCommon/VectorPattern.h b/include/mlir/Conversion/LLVMCommon/VectorPattern.h index 65988a2..fdc787d 100644 --- a/include/mlir/Conversion/LLVMCommon/VectorPattern.h +++ b/include/mlir/Conversion/LLVMCommon/VectorPattern.h
@@ -70,13 +70,16 @@ 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/include/mlir/Dialect/GPU/IR/GPUOps.td b/include/mlir/Dialect/GPU/IR/GPUOps.td index 13fe6d3..e6573c9 100644 --- a/include/mlir/Dialect/GPU/IR/GPUOps.td +++ b/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -450,7 +450,7 @@ bool isKernel() { if (getKernel()) return true; - return (*this)->getAttrOfType<UnitAttr>( + return (*this)->getDiscardableAttrOfType<UnitAttr>( GPUDialect::getKernelFuncAttrName()) != nullptr; }
diff --git a/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td b/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td index 2834ae1..10e4270 100644 --- a/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td +++ b/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td
@@ -157,13 +157,16 @@ 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/include/mlir/Dialect/LLVMIR/LLVMInterfaces.td b/include/mlir/Dialect/LLVMIR/LLVMInterfaces.td index c59b38c..66a1fe0 100644 --- a/include/mlir/Dialect/LLVMIR/LLVMInterfaces.td +++ b/include/mlir/Dialect/LLVMIR/LLVMInterfaces.td
@@ -35,6 +35,17 @@ 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/include/mlir/Dialect/LLVMIR/NVVMOps.td b/include/mlir/Dialect/LLVMIR/NVVMOps.td index 634918b..9671dca 100644 --- a/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ b/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -279,7 +279,8 @@ 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/include/mlir/Dialect/XeGPU/Transforms/XeGPULayoutImpl.h b/include/mlir/Dialect/XeGPU/Transforms/XeGPULayoutImpl.h index defc389..9d0ee04 100644 --- a/include/mlir/Dialect/XeGPU/Transforms/XeGPULayoutImpl.h +++ b/include/mlir/Dialect/XeGPU/Transforms/XeGPULayoutImpl.h
@@ -93,6 +93,10 @@ /// 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/include/mlir/IR/SymbolInterfaces.td b/include/mlir/IR/SymbolInterfaces.td index 1a6190c..97570b1 100644 --- a/include/mlir/IR/SymbolInterfaces.td +++ b/include/mlir/IR/SymbolInterfaces.td
@@ -34,7 +34,12 @@ "::mlir::StringAttr", "getNameAttr", (ins) >, InterfaceMethod<"Sets the name of this symbol.", - "void", "setName", (ins "::mlir::StringAttr":$name) + "void", "setSymbolName", (ins "::mlir::StringAttr":$name), [{}], + /*defaultImplementation=*/[{ + auto *op = this->getOperation(); + op->setInherentAttr(::mlir::StringAttr::get( + op->getContext(), mlir::SymbolTable::getSymbolAttrName()), name); + }] >, InterfaceMethod<"Gets the visibility of this symbol.", "mlir::SymbolTable::Visibility", "getVisibility", (ins) @@ -135,9 +140,19 @@ 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.
diff --git a/include/mlir/IR/SymbolTable.h b/include/mlir/IR/SymbolTable.h index 96be32e..a6d8e86 100644 --- a/include/mlir/IR/SymbolTable.h +++ b/include/mlir/IR/SymbolTable.h
@@ -118,6 +118,9 @@ /// 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"; } /// Sets the name of the given symbol operation. static void setSymbolName(Operation *symbol, StringAttr name);
diff --git a/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp b/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp index 2a84633..2f6ede5 100644 --- a/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp +++ b/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp
@@ -3218,9 +3218,9 @@ // 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/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp b/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp index d164f0e..982b1e5 100644 --- a/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp +++ b/lib/Conversion/ArithToLLVM/ArithToLLVM.cpp
@@ -114,9 +114,10 @@ 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 = @@ -203,7 +204,7 @@ 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/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp b/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp index ceda388..74acbe3 100644 --- a/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp +++ b/lib/Conversion/ComplexToLLVM/ComplexToLLVM.cpp
@@ -96,8 +96,8 @@ 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/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp b/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp index fef78a4..2c2b6b5 100644 --- a/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp +++ b/lib/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.cpp
@@ -133,6 +133,18 @@ return result; } +/// 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 + discardableAttrs.push_back(attr); + } + op->setDiscardableAttrs(discardableAttrs); +} + /// Convert the destination block signature (if necessary) and lower the branch /// op to llvm.br. struct BranchOpLowering : public ConvertOpToLLVMPattern<cf::BranchOp> { @@ -148,12 +160,12 @@ 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 +203,7 @@ *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/lib/Conversion/GPUCommon/GPUOpsLowering.cpp b/lib/Conversion/GPUCommon/GPUOpsLowering.cpp index e2d251b..66322f9 100644 --- a/lib/Conversion/GPUCommon/GPUOpsLowering.cpp +++ b/lib/Conversion/GPUCommon/GPUOpsLowering.cpp
@@ -630,7 +630,6 @@ 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 @@ 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 @@ // 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 @@ 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/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h b/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h index 777f3b9..2a6036a 100644 --- a/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h +++ b/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
@@ -87,7 +87,7 @@ 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/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp b/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp index 80420c2..2b57e90 100644 --- a/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp +++ b/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
@@ -513,8 +513,8 @@ // 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/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp index a3819df..ab7666b 100644 --- a/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp +++ b/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
@@ -302,7 +302,7 @@ 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 @@ 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/lib/Conversion/LLVMCommon/LowerFunctionDiscardablesToLLVM.cpp b/lib/Conversion/LLVMCommon/LowerFunctionDiscardablesToLLVM.cpp index 5427877..859d96d 100644 --- a/lib/Conversion/LLVMCommon/LowerFunctionDiscardablesToLLVM.cpp +++ b/lib/Conversion/LLVMCommon/LowerFunctionDiscardablesToLLVM.cpp
@@ -29,7 +29,8 @@ 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/lib/Conversion/LLVMCommon/Pattern.cpp b/lib/Conversion/LLVMCommon/Pattern.cpp index 9c116ef..2c4b8db 100644 --- a/lib/Conversion/LLVMCommon/Pattern.cpp +++ b/lib/Conversion/LLVMCommon/Pattern.cpp
@@ -372,7 +372,19 @@ auto callIntrOp = LLVM::CallIntrinsicOp::create( rewriter, loc, resType, rewriter.getStringAttr(intrinsic), operands); // Propagate attributes. - callIntrOp->setAttrs(op->getAttrDictionary()); + SmallVector<NamedAttribute> discardableAttrs; + auto copyAttr = [&](StringAttr name, Attribute attr) { + if (callIntrOp->getInherentAttr(name).has_value()) + callIntrOp->setInherentAttr(name, attr); + else + 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/lib/Conversion/MathToXeVM/MathToXeVM.cpp b/lib/Conversion/MathToXeVM/MathToXeVM.cpp index 7143dcc..58b7948 100644 --- a/lib/Conversion/MathToXeVM/MathToXeVM.cpp +++ b/lib/Conversion/MathToXeVM/MathToXeVM.cpp
@@ -144,8 +144,16 @@ // 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->setAttr(fastAttr.getName(), 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/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp b/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp index 7207c96..4bc0d3a 100644 --- a/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp +++ b/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
@@ -353,7 +353,7 @@ 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 @@ "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/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp b/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp index d6ff850..662f83c 100644 --- a/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp +++ b/lib/Conversion/XeGPUToXeVM/XeGPUToXeVM.cpp
@@ -804,7 +804,7 @@ 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 @@ 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/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp b/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp index d2594ca..e653c27 100644 --- a/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp +++ b/lib/Conversion/XeVMToLLVM/XeVMToLLVM.cpp
@@ -259,8 +259,9 @@ } 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()); @@ -269,8 +270,9 @@ } 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()); @@ -551,7 +553,20 @@ funcOp.setArgAttr(idx, attrName, rewriter.getUnitAttr()); auto callOp = LLVM::CallOp::create(rewriter, loc, funcOp, args); - callOp->setAttrs(funcOp->getAttrs()); + SmallVector<NamedAttribute> discardableAttrs; + auto copyAttr = [&](StringAttr name, Attribute attr) { + if (callOp->getInherentAttr(name).has_value()) + callOp->setInherentAttr(name, attr); + else + discardableAttrs.emplace_back(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); + }); + callOp->setDiscardableAttrs(discardableAttrs); return callOp; } @@ -986,14 +1001,15 @@ 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(); } @@ -1009,7 +1025,7 @@ // 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(); } @@ -1947,7 +1963,7 @@ 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/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp b/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp index 332ac8c..b59f7e9 100644 --- a/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp +++ b/lib/Dialect/AMDGPU/Transforms/EmulateAtomics.cpp
@@ -110,7 +110,10 @@ 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/lib/Dialect/AMDGPU/Transforms/MaskedloadToLoad.cpp b/lib/Dialect/AMDGPU/Transforms/MaskedloadToLoad.cpp index 89ef51f..8c2d697 100644 --- a/lib/Dialect/AMDGPU/Transforms/MaskedloadToLoad.cpp +++ b/lib/Dialect/AMDGPU/Transforms/MaskedloadToLoad.cpp
@@ -81,7 +81,7 @@ 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 @@ 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/lib/Dialect/GPU/IR/GPUDialect.cpp b/lib/Dialect/GPU/IR/GPUDialect.cpp index f385522..f414596 100644 --- a/lib/Dialect/GPU/IR/GPUDialect.cpp +++ b/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -266,7 +266,7 @@ 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 @@ 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 @@ 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 @@ 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 @@ 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 @@ // 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 @@ 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 @@ 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 @@ 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 @@ } LogicalResult GPUModuleOp::verify() { - auto targets = getOperation()->getAttrOfType<ArrayAttr>("targets"); + auto targets = getTargetsAttr(); if (!targets) return success(); @@ -2453,8 +2461,7 @@ 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 @@ 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/lib/Dialect/GPU/IR/InferIntRangeInterfaceImpls.cpp b/lib/Dialect/GPU/IR/InferIntRangeInterfaceImpls.cpp index 3369583..b15deed 100644 --- a/lib/Dialect/GPU/IR/InferIntRangeInterfaceImpls.cpp +++ b/lib/Dialect/GPU/IR/InferIntRangeInterfaceImpls.cpp
@@ -68,7 +68,9 @@ 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/lib/Dialect/GPU/Transforms/EliminateBarriers.cpp b/lib/Dialect/GPU/Transforms/EliminateBarriers.cpp index 2116254..05ede7b 100644 --- a/lib/Dialect/GPU/Transforms/EliminateBarriers.cpp +++ b/lib/Dialect/GPU/Transforms/EliminateBarriers.cpp
@@ -45,7 +45,7 @@ /// 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/lib/Dialect/GPU/Transforms/KernelOutlining.cpp b/lib/Dialect/GPU/Transforms/KernelOutlining.cpp index 831ac7f..0ee3d42 100644 --- a/lib/Dialect/GPU/Transforms/KernelOutlining.cpp +++ b/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
@@ -394,8 +394,9 @@ // 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 @@ // 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/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp b/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp index 594c7a2..07b6dd4 100644 --- a/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp +++ b/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp
@@ -44,8 +44,9 @@ 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 @@ 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/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp b/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp index b81b312..317d6b6 100644 --- a/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp +++ b/lib/Dialect/LLVMIR/IR/LLVMAttrs.cpp
@@ -620,7 +620,21 @@ auto parentFunction = op->getParentOfType<FunctionOpInterface>(); if (!parentFunction) return {}; - return parentFunction.getOperation()->getAttrOfType<TargetFeaturesAttr>( + 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()); }
diff --git a/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp b/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp index 1e1ec23..3758a11 100644 --- a/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp +++ b/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp
@@ -50,6 +50,13 @@ 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 attrs; +} + static auto processFMFAttr(ArrayRef<NamedAttribute> attrs) { SmallVector<NamedAttribute, 8> filteredAttrs( llvm::make_filter_range(attrs, [&](NamedAttribute attr) { @@ -278,14 +285,15 @@ 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 @@ 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 @@ getOpBundleOperands().getTypes(), getOpBundleTags()); } - p.printOptionalAttrDict(processFMFAttr((*this)->getAttrs()), + p.printOptionalAttrDict(processFMFAttr(getAttrsForPrinting(*this).getAttrs()), {getCalleeAttrName(), getTailCallKindAttrName(), getVarCalleeTypeAttrName(), getCConvAttrName(), getOperandSegmentSizesAttrName(), @@ -1707,7 +1712,7 @@ getOpBundleOperands().getTypes(), getOpBundleTags()); } - p.printOptionalAttrDict((*this)->getAttrs(), + p.printOptionalAttrDict(getAttrsForPrinting(*this).getAttrs(), {getCalleeAttrName(), getOperandSegmentSizeAttr(), getCConvAttrName(), getVarCalleeTypeAttrName(), getOpBundleSizesAttrName(), @@ -1887,7 +1892,7 @@ << value.getType() << ") "; } - p.printOptionalAttrDict((*this)->getAttrs(), {"cleanup"}); + p.printOptionalAttrDict(getAttrsForPrinting(*this).getAttrs(), {"cleanup"}); p << ": " << getType(); } @@ -4259,7 +4264,7 @@ getOpBundleOperands().getTypes(), getOpBundleTagsAttr()); } - p.printOptionalAttrDict(processFMFAttr((*this)->getAttrs()), + p.printOptionalAttrDict(processFMFAttr(getAttrsForPrinting(*this).getAttrs()), {getOperandSegmentSizesAttrName(), getOpBundleSizesAttrName(), getIntrinAttrName(), getOpBundleTagsAttrName(), getArgAttrsAttrName(),
diff --git a/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp index 0a9ba99..8433447 100644 --- a/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp +++ b/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -1082,7 +1082,8 @@ getLayoutBAttrName(), getMultiplicandAPtxTypeAttrName(), getMultiplicandBPtxTypeAttrName()}); - p.printOptionalAttrDict(this->getOperation()->getAttrs(), ignoreAttrNames); + p.printOptionalAttrDict((*this)->getDiscardableAttrDictionary().getValue(), + ignoreAttrNames); // Print the types of the operands and result. p << " : " << "("; @@ -1575,7 +1576,8 @@ getMultiplicandAPtxTypeAttrName(), getMultiplicandBPtxTypeAttrName(), getOrderedMetadataAttrName(), getKindAttrName()}); - p.printOptionalAttrDict((*this)->getAttrs(), ignoreAttrNames); + p.printOptionalAttrDict((*this)->getDiscardableAttrDictionary().getValue(), + ignoreAttrNames); p << " : "; p << "("; for (int i = 0; i < 3; ++i) { @@ -2165,7 +2167,8 @@ getMultiplicandBPtxTypeAttrName(), getScaleVecSizeAttrName(), getBlockScaleFormatAttrName(), getKindAttrName()}); - p.printOptionalAttrDict(this->getOperation()->getAttrs(), ignoreAttrNames); + p.printOptionalAttrDict((*this)->getDiscardableAttrDictionary().getValue(), + ignoreAttrNames); // Print type signature p << " : ("; @@ -2423,7 +2426,8 @@ getOrderedMetadataAttrName(), getScaleVecSizeAttrName(), getBlockScaleFormatAttrName(), getKindAttrName()}); - p.printOptionalAttrDict(this->getOperation()->getAttrs(), ignoreAttrNames); + p.printOptionalAttrDict((*this)->getDiscardableAttrDictionary().getValue(), + ignoreAttrNames); // Print type signature p << " : ("; @@ -5817,12 +5821,13 @@ /// 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()); } @@ -7006,8 +7011,8 @@ } // 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 " << "'" @@ -7026,7 +7031,7 @@ 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/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp b/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp index 657835a..9409141 100644 --- a/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp +++ b/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
@@ -286,7 +286,7 @@ continue; ssaCopies.insert(ssaCopy); - if (!ssaCopy->hasAttr(LLVM::LLVMDialect::getNoAliasAttrName())) + if (!ssaCopy->hasDiscardableAttr(LLVM::LLVMDialect::getNoAliasAttrName())) continue; noAliasParams.insert(ssaCopy); }
diff --git a/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp b/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp index a5fccbd..10d203c 100644 --- a/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp +++ b/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp
@@ -25,8 +25,9 @@ : 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/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp b/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp index 642e588..2649564 100644 --- a/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp +++ b/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
@@ -272,7 +272,7 @@ return verifyMmaSyncOp(this->getOperation(), getMatrixA(), getMatrixB(), getMatrixC(), getMmaShapeAsArray(), - getOperation()->hasAttr(getTf32EnabledAttrName())); + getTf32Enabled().value_or(false)); } //===----------------------------------------------------------------------===// @@ -296,8 +296,7 @@ return verifyMmaSyncOp(this->getOperation(), getMatrixA(), getMatrixB(), getMatrixC(), getMmaShapeAsArray(), - getOperation()->hasAttr(getTf32EnabledAttrName()), - true); + getTf32Enabled().value_or(false), true); } //===----------------------------------------------------------------------===//
diff --git a/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp b/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp index 7f626a6..118796e 100644 --- a/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp +++ b/lib/Dialect/NVGPU/Transforms/MmaSyncTF32Transform.cpp
@@ -35,7 +35,7 @@ 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/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp b/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp index 9d16864..d9648a6 100644 --- a/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp +++ b/lib/Dialect/XeGPU/Transforms/XeGPUBlocking.cpp
@@ -433,14 +433,15 @@ 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 @@ // 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,10 @@ // 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); + xegpu::dropInstDataOnInherentAttrs(op); }); // Resolve UnrealizedConversionCastOps generated by SCF structural type
diff --git a/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp b/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp index 076b937..f696020 100644 --- a/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp +++ b/lib/Dialect/XeGPU/Transforms/XeGPULayoutImpl.cpp
@@ -70,6 +70,13 @@ 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) { @@ -391,8 +398,8 @@ 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 @@ -407,19 +414,19 @@ 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/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp b/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp index a7534d3..2903e19 100644 --- a/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp +++ b/lib/Dialect/XeGPU/Transforms/XeGPUPropagateLayout.cpp
@@ -1752,10 +1752,11 @@ conflictingCreateNdOp.getContext(), currTDescType.getShape(), currTDescType.getElementType(), currTDescType.getEncoding(), expectedLayout); - xegpu::CreateNdDescOp newOp = xegpu::CreateNdDescOp::create( - builder, consumerOp->getLoc(), newTensorDescType, + auto newOp = xegpu::CreateNdDescOp::create( + builder, consumerOp->getLoc(), TypeRange{newTensorDescType}, conflictingCreateNdOp->getOperands(), - conflictingCreateNdOp->getAttrs()); + conflictingCreateNdOp.getProperties(), + conflictingCreateNdOp->getDiscardableAttrDictionary().getValue()); // Replace the tensor descriptor operand in the consumer op with the new // tensor descriptor. consumerOp->replaceUsesOfWith(tdescValue, newOp.getResult());
diff --git a/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp b/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp index bc85140..8122330 100644 --- a/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp +++ b/lib/Dialect/XeGPU/Transforms/XeGPUSgToLaneDistribute.cpp
@@ -128,8 +128,9 @@ return failure(); auto newOp = xegpu::CreateNdDescOp::create( - rewriter, op.getLoc(), resultType.dropLayouts(), op.getOperands(), - op->getAttrs()); + rewriter, op.getLoc(), TypeRange{resultType.dropLayouts()}, + op.getOperands(), op.getProperties(), + op->getDiscardableAttrDictionary().getValue()); rewriter.replaceOp(op, newOp.getResult()); return success(); } @@ -345,10 +346,11 @@ 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/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp b/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp index 5914e23..1854c63 100644 --- a/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp +++ b/lib/Dialect/XeGPU/Transforms/XeGPUUnroll.cpp
@@ -428,9 +428,12 @@ if (tmpC) operands.push_back(tmpC); - tmpC = xegpu::DpasOp::create( - rewriter, loc, vecTy, operands, - xegpu::dropInstDataOnAttrs(op->getAttrs())); + auto newDpasOp = xegpu::DpasOp::create( + rewriter, loc, TypeRange{vecTy}, operands, op.getProperties(), + xegpu::dropInstDataOnAttrs( + op->getDiscardableAttrDictionary().getValue())); + xegpu::dropInstDataOnInherentAttrs(newDpasOp); + tmpC = newDpasOp.getResult(); } newOps.push_back(tmpC); } @@ -531,8 +534,10 @@ bScaleVals[batch * (kIters * nIters) + k * nIters + j]); newDpasMxOp = xegpu::DpasMxOp::create( - rewriter, loc, vecTy, operands, - xegpu::dropInstDataOnAttrs(op->getAttrs())); + rewriter, loc, TypeRange{vecTy}, operands, op.getProperties(), + xegpu::dropInstDataOnAttrs( + op->getDiscardableAttrDictionary().getValue())); + xegpu::dropInstDataOnInherentAttrs(newDpasMxOp); tmpC = newDpasMxOp.getResult(); } newOps.push_back(newDpasMxOp); @@ -563,11 +568,7 @@ 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 +658,7 @@ 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/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp b/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp index 0e8a386..accc401 100644 --- a/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp +++ b/lib/Dialect/XeGPU/Transforms/XeGPUWgToSgDistribute.cpp
@@ -41,7 +41,7 @@ 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 @@ 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/lib/Dialect/XeGPU/Utils/XeGPUUtils.cpp b/lib/Dialect/XeGPU/Utils/XeGPUUtils.cpp index f351d91..bd2fe1e 100644 --- a/lib/Dialect/XeGPU/Utils/XeGPUUtils.cpp +++ b/lib/Dialect/XeGPU/Utils/XeGPUUtils.cpp
@@ -161,9 +161,10 @@ } 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; } } @@ -265,8 +266,9 @@ } 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; } @@ -323,11 +325,11 @@ } 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); } } @@ -371,11 +373,11 @@ } 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); } } @@ -385,8 +387,9 @@ 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; } @@ -403,11 +406,11 @@ 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/lib/IR/SymbolTable.cpp b/lib/IR/SymbolTable.cpp index 31e9ae5..b028818 100644 --- a/lib/IR/SymbolTable.cpp +++ b/lib/IR/SymbolTable.cpp
@@ -295,7 +295,7 @@ /// 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/lib/Target/LLVM/XeVM/Target.cpp b/lib/Target/LLVM/XeVM/Target.cpp index 18f6b7b..731ac4b 100644 --- a/lib/Target/LLVM/XeVM/Target.cpp +++ b/lib/Target/LLVM/XeVM/Target.cpp
@@ -619,7 +619,8 @@ 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/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMIRToLLVMTranslation.cpp b/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMIRToLLVMTranslation.cpp index b6dee4e..6733d6e 100644 --- a/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMIRToLLVMTranslation.cpp +++ b/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMIRToLLVMTranslation.cpp
@@ -285,7 +285,7 @@ } mlirMmra = ArrayAttr::get(ctx, tags); } - op->setAttr(LLVMDialect::getMmraAttrName(), mlirMmra); + op->setDiscardableAttr(LLVMDialect::getMmraAttrName(), mlirMmra); return success(); }
diff --git a/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp b/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp index 3e0a228..aa62d0d 100644 --- a/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp +++ b/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp
@@ -639,7 +639,7 @@ 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/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp b/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp index 4239321..18eb07a 100644 --- a/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp +++ b/lib/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.cpp
@@ -195,8 +195,8 @@ 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 @@ 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/lib/Target/LLVMIR/ModuleImport.cpp b/lib/Target/LLVMIR/ModuleImport.cpp index ef2363f..bb85fe8 100644 --- a/lib/Target/LLVMIR/ModuleImport.cpp +++ b/lib/Target/LLVMIR/ModuleImport.cpp
@@ -1039,8 +1039,9 @@ 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(); } @@ -1056,8 +1057,9 @@ 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(); } @@ -1073,8 +1075,9 @@ 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(); } @@ -1194,13 +1197,13 @@ 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())); } @@ -1216,8 +1219,8 @@ 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() { @@ -1299,7 +1302,7 @@ 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/lib/Target/LLVMIR/ModuleTranslation.cpp b/lib/Target/LLVMIR/ModuleTranslation.cpp index 76fa43c..26a497d 100644 --- a/lib/Target/LLVMIR/ModuleTranslation.cpp +++ b/lib/Target/LLVMIR/ModuleTranslation.cpp
@@ -950,9 +950,11 @@ 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(); @@ -983,7 +985,7 @@ 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(); @@ -2500,7 +2502,7 @@ } 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(); @@ -2514,7 +2516,7 @@ } 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/lib/Target/LLVMIR/Transforms/TargetToDataLayout.cpp b/lib/Target/LLVMIR/Transforms/TargetToDataLayout.cpp index c0f9ceb..1246c9c 100644 --- a/lib/Target/LLVMIR/Transforms/TargetToDataLayout.cpp +++ b/lib/Target/LLVMIR/Transforms/TargetToDataLayout.cpp
@@ -33,7 +33,7 @@ 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 @@ 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/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp b/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp index 24f37cc..0cc5f5a 100644 --- a/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp +++ b/lib/Target/LLVMIR/Transforms/TargetToTargetFeatures.cpp
@@ -37,7 +37,7 @@ 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 @@ LLVM::TargetAttr::get(&getContext(), targetAttr.getTriple(), targetAttr.getChip(), fullTargetFeaturesAttr); - op->setAttr(LLVM::LLVMDialect::getTargetAttrName(), updatedTargetAttr); + op->setDiscardableAttr(LLVM::LLVMDialect::getTargetAttrName(), + updatedTargetAttr); } };
diff --git a/test/Conversion/ArithToLLVM/attribute-storage.mlir b/test/Conversion/ArithToLLVM/attribute-storage.mlir new file mode 100644 index 0000000..5ccc0f0 --- /dev/null +++ b/test/Conversion/ArithToLLVM/attribute-storage.mlir
@@ -0,0 +1,40 @@ +// 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 + + // 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/test/Conversion/MathToXeVM/math-to-xevm.mlir b/test/Conversion/MathToXeVM/math-to-xevm.mlir index 0ca8d14..55eb300 100644 --- a/test/Conversion/MathToXeVM/math-to-xevm.mlir +++ b/test/Conversion/MathToXeVM/math-to-xevm.mlir
@@ -67,6 +67,9 @@ // 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/test/Dialect/XeGPU/xegpu-blocking.mlir b/test/Dialect/XeGPU/xegpu-blocking.mlir index 44b7615..c7cbdf1 100644 --- a/test/Dialect/XeGPU/xegpu-blocking.mlir +++ b/test/Dialect/XeGPU/xegpu-blocking.mlir
@@ -26,7 +26,7 @@ %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 @@ %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 @@ %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 @@ %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 @@ %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 @@ %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/test/lib/Dialect/LLVM/TestPatterns.cpp b/test/lib/Dialect/LLVM/TestPatterns.cpp index 69a3d98..2252db9 100644 --- a/test/lib/Dialect/LLVM/TestPatterns.cpp +++ b/test/lib/Dialect/LLVM/TestPatterns.cpp
@@ -79,8 +79,9 @@ 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/test/lib/Dialect/Test/TestOpDefs.cpp b/test/lib/Dialect/Test/TestOpDefs.cpp index 06871fd..4c68335 100644 --- a/test/lib/Dialect/Test/TestOpDefs.cpp +++ b/test/lib/Dialect/Test/TestOpDefs.cpp
@@ -30,10 +30,6 @@ return getSymNameAttr(); } -void OverriddenSymbolVisibilityOp::setName(StringAttr name) { - setSymNameAttr(name); -} - static StringLiteral getVisibilityString(SymbolTable::Visibility visibility) { switch (visibility) { case SymbolTable::Visibility::Private:
diff --git a/test/lib/Dialect/XeGPU/TestXeGPUTransforms.cpp b/test/lib/Dialect/XeGPU/TestXeGPUTransforms.cpp index f6b0c50..a481352 100644 --- a/test/lib/Dialect/XeGPU/TestXeGPUTransforms.cpp +++ b/test/lib/Dialect/XeGPU/TestXeGPUTransforms.cpp
@@ -166,7 +166,7 @@ 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 @@ 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/unittests/Dialect/LLVMIR/CMakeLists.txt b/unittests/Dialect/LLVMIR/CMakeLists.txt index 7cc130d..dc0e82d 100644 --- a/unittests/Dialect/LLVMIR/CMakeLists.txt +++ b/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/unittests/Dialect/LLVMIR/LLVMAttrsTest.cpp b/unittests/Dialect/LLVMIR/LLVMAttrsTest.cpp new file mode 100644 index 0000000..15b5464 --- /dev/null +++ b/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/unittests/IR/SymbolTableTest.cpp b/unittests/IR/SymbolTableTest.cpp index f18cf85..fd96467 100644 --- a/unittests/IR/SymbolTableTest.cpp +++ b/unittests/IR/SymbolTableTest.cpp
@@ -89,7 +89,7 @@ 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());
diff --git a/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp b/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp index 5271923..c2ca299 100644 --- a/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp +++ b/unittests/Target/LLVM/SerializeToLLVMBitcode.cpp
@@ -105,7 +105,8 @@ 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 @@ 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.