[mlir] Remove types from attributes

[Mogball]

Summary

This patch removes the type field from Attribute along with the
Attribute::getType accessor.

Going forward, this means that attributes in MLIR will no longer have
types as a first-class concept. This patch lays the groundwork to
incrementally remove or refactor code that relies on generic attributes
being typed. The immediate impact will be on attributes that rely on
Attribute containing a type, such as IntegerAttr,
DenseElementsAttr, and ml_program::ExternAttr, which will now need
to define a type parameter on their storage classes. This will save
memory as all other attribute kinds will no longer contain a type.

Moreover, it will not be possible to generically query the type of an
attribute directly. This patch provides an attribute interface
TypedAttr that implements only one method, getType, which can be
used to generically query the types of attributes that implement the
interface. This interface can be used to retain the concept of a "typed
attribute". The ODS-generated accessor for a type parameter
automatically implements this method.

Next steps will be to refactor the assembly formats of certain operations
that rely on parseAttribute(type) and printAttributeWithoutType to
remove special handling of type elision until type can be removed from
the dialect parsing hook entirely; and incrementally remove uses of
TypedAttr.

[mehdi_amini]

mlir/include/mlir/IR/AttrTypeBase.td

391

Can you document the typeBuilder parameter as well as this logic?

[lattner]

Oh nice, TypedAttr is a great idea. I'm very +1 on this direction. Thank you Jeff!

mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp

2415

If the types are the same, the attributes should unique to be the same thing, can you just change this to something like:
if (re != im || !re) ?

mlir/lib/Parser/DialectSymbolParser.cpp

253

Super nit but you can use:

if (auto typedAttr = attr.dyn_cast_or_null<TypedAttr>()) {

[rriddle]

No doc updates were necessary?

mlir/include/mlir/IR/BuiltinAttributeInterfaces.td

160

483–484

[Mogball]

Review comments and update docs a little.

mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp

2415

TypedAttr is an interface so the underlying attributes can still be different, e.g. [-1.0 : f32, 1.0 : f32] is valid

[lattner]

Oh right, of course, thx.

[jpienaar]

Nice, thanks!

mlir/include/mlir/IR/BuiltinAttributeInterfaces.td

478

With a type?

mlir/include/mlir/IR/BuiltinAttributes.td

200–202

Could you add what the shaped type represents?

mlir/lib/IR/BuiltinAttributes.cpp

690–692

OOC why is this called tblgenKey?

988

Shouldn't this be in the attributes verify somewhere?

[jloser]

mlir/docs/AttributesAndTypes.md

716 ↗

(On Diff #445891)

Super nit: s/Parameterk/Parameter

[Mogball]

mlir/lib/IR/BuiltinAttributes.cpp

988

This might be legacy code. There are a lot of redundant asserts around here

[Mogball]

review comments

[Mogball]

rebase

[lattner]

huzzah, congrats Jeff!

[dcaballe]

QQ: Should a TypeAttr also implements the TypeAttrInterface. The fact that is not is causing some problems with IREE integration.

[mehdi_amini]

I don’t get why it should? What would be the type of a type attribute? Returning the holder type does not seems consistent to me: I see TypedAttr interface to indicate the type of the data wrapped by the attribute, we would need some sort of meta-type system.
What is the use case for IREE?

[dcaballe]

I see your point. It makes sense to me. It looks like we had some code that was resilient to null types and I assumed operands should be converted to TypedAttr, which doesn't seem to be the case. Thanks for clarifying!

mlir/include/mlir/Dialect/Arithmetic/IR/ArithmeticOps.td

 #define ARITHMETIC_OPS
 
 include "mlir/Dialect/Arithmetic/IR/ArithmeticBase.td"
 include "mlir/Interfaces/CastInterfaces.td"
 include "mlir/Interfaces/InferIntRangeInterface.td"
 include "mlir/Interfaces/InferTypeOpInterface.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 include "mlir/Interfaces/VectorInterfaces.td"
+include "mlir/IR/BuiltinAttributeInterfaces.td"
 include "mlir/IR/OpAsmInterface.td"
 
 // Base class for Arithmetic dialect ops. Ops in this dialect have no side
 // effects and can be applied element-wise to vectors and tensors.
 class Arith_Op<string mnemonic, list<Trait> traits = []> :
     Op<Arithmetic_Dialect, mnemonic, traits # [NoSideEffect,
     DeclareOpInterfaceMethods<VectorUnrollOpInterface>] #
     ElementwiseMappable.traits>;
let description = [{
     // Integer constant
     %1 = arith.constant 42 : i32
 
     // Equivalent generic form
     %1 = "arith.constant"() {value = 42 : i32} : () -> i32
     ```
   }];
 
-  let arguments = (ins AnyAttr:$value);
+  let arguments = (ins TypedAttrInterface:$value);
   // TODO: Disallow arith.constant to return anything other than a signless
   // integer or float like. Downstream users of Arithmetic should only be
   // working with signless integers, floats, or vectors/tensors thereof.
   // However, it is necessary to allow arith.constant to return vectors/tensors
   // of strings and signed/unsigned integers (for now) as an artefact of
   // splitting the Standard dialect.
   let results = (outs /*SignlessIntegerOrFloatLike*/AnyType:$result);
 

mlir/include/mlir/Dialect/CommonFolders.h

template <
     class AttrElementT, class ElementValueT = typename AttrElementT::ValueType,
     class CalculationT =
         function_ref<Optional<ElementValueT>(ElementValueT, ElementValueT)>>
 Attribute constFoldBinaryOpConditional(ArrayRef<Attribute> operands,
                                        const CalculationT &calculate) {
   assert(operands.size() == 2 && "binary op takes two operands");
   if (!operands[0] || !operands[1])
     return {};
-  if (operands[0].getType() != operands[1].getType())
-    return {};
 
   if (operands[0].isa<AttrElementT>() && operands[1].isa<AttrElementT>()) {
     auto lhs = operands[0].cast<AttrElementT>();
     auto rhs = operands[1].cast<AttrElementT>();
+    if (lhs.getType() != rhs.getType())
+      return {};
 
     auto calRes = calculate(lhs.getValue(), rhs.getValue());
 
     if (!calRes)
       return {};
 
     return AttrElementT::get(lhs.getType(), *calRes);
   }
 
   if (operands[0].isa<SplatElementsAttr>() &&
       operands[1].isa<SplatElementsAttr>()) {
     // Both operands are splats so we can avoid expanding the values out and
     // just fold based on the splat value.
     auto lhs = operands[0].cast<SplatElementsAttr>();
     auto rhs = operands[1].cast<SplatElementsAttr>();
+    if (lhs.getType() != rhs.getType())
+      return {};
 
     auto elementResult = calculate(lhs.getSplatValue<ElementValueT>(),
                                    rhs.getSplatValue<ElementValueT>());
     if (!elementResult)
       return {};
 
     return DenseElementsAttr::get(lhs.getType(), *elementResult);
   } else if (operands[0].isa<ElementsAttr>() &&
              operands[1].isa<ElementsAttr>()) {
     // Operands are ElementsAttr-derived; perform an element-wise fold by
     // expanding the values.
     auto lhs = operands[0].cast<ElementsAttr>();
     auto rhs = operands[1].cast<ElementsAttr>();
+    if (lhs.getType() != rhs.getType())
+      return {};
 
     auto lhsIt = lhs.value_begin<ElementValueT>();
     auto rhsIt = rhs.value_begin<ElementValueT>();
     SmallVector<ElementValueT, 4> elementResults;
     elementResults.reserve(lhs.getNumElements());
     for (size_t i = 0, e = lhs.getNumElements(); i < e; ++i, ++lhsIt, ++rhsIt) {
       auto elementResult = calculate(*lhsIt, *rhsIt);
       if (!elementResult)

mlir/include/mlir/Dialect/EmitC/IR/EmitC.td

let description = [{
 
     // Constant emitted as `char = CHAR_MIN;`
     %1 = "emitc.constant"()
         {value = #emitc.opaque<"CHAR_MIN"> : !emitc.opaque<"char">}
         : () -> !emitc.opaque<"char">
     ```
   }];
 
-  let arguments = (ins AnyAttr:$value);
+  let arguments = (ins TypedAttrInterface:$value);
   let results = (outs AnyType);
 
   let hasFolder = 1;
   // TODO: Disallow empty constants.
   let hasVerifier = 1;
 }
 
 def EmitC_IncludeOp
let description = [{
     %0 = "emitc.variable"() {value = 0 : i32} : () -> i32
     %1 = "emitc.variable"() {value = 0 : i32} : () -> i32
     %2 = emitc.apply "&"(%0) : (i32) -> !emitc.ptr<i32>
     %3 = emitc.apply "&"(%1) : (i32) -> !emitc.ptr<i32>
     emitc.call "write"(%2, %3) : (!emitc.ptr<i32>, !emitc.ptr<i32>) -> ()
     ```
   }];
 
-  let arguments = (ins AnyAttr:$value);
+  let arguments = (ins TypedAttrInterface:$value);
   let results = (outs AnyType);
 
   let hasVerifier = 1;
 }
 
 #endif // MLIR_DIALECT_EMITC_IR_EMITC

mlir/include/mlir/Dialect/EmitC/IR/EmitCAttributes.td

 //===- EmitCAttributes.td - EmitC attributes ---------------*- tablegen -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Defines the MLIR EmitC attributes.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_EMITC_IR_EMITCATTRIBUTES
 #define MLIR_DIALECT_EMITC_IR_EMITCATTRIBUTES
 
 include "mlir/IR/AttrTypeBase.td"
+include "mlir/IR/BuiltinAttributeInterfaces.td"
 include "mlir/Dialect/EmitC/IR/EmitCBase.td"
 
 //===----------------------------------------------------------------------===//
 // EmitC attribute definitions
 //===----------------------------------------------------------------------===//
 
-class EmitC_Attr<string name, string attrMnemonic>
-    : AttrDef<EmitC_Dialect, name> {
+class EmitC_Attr<string name, string attrMnemonic, list<Trait> traits = []>
+    : AttrDef<EmitC_Dialect, name, traits> {
   let mnemonic = attrMnemonic;
 }
 
-def EmitC_OpaqueAttr : EmitC_Attr<"Opaque", "opaque"> {
+def EmitC_OpaqueAttr : EmitC_Attr<"Opaque", "opaque", [TypedAttrInterface]> {
   let summary = "An opaque attribute";
 
   let description = [{
     An opaque attribute of which the value gets emitted as is.
 
     Example:
 
     ```mlir
     #emitc.opaque<"">
     #emitc.opaque<"NULL">
     #emitc.opaque<"nullptr">
     ```
   }];
 
-  let parameters = (ins StringRefParameter<"the opaque value">:$value);
-  
+  let parameters = (ins "Type":$type,
+                        StringRefParameter<"the opaque value">:$value);
   let hasCustomAssemblyFormat = 1;
 }
 
 #endif // MLIR_DIALECT_EMITC_IR_EMITCATTRIBUTES

mlir/include/mlir/Dialect/MLProgram/IR/MLProgramAttributes.h

 //===- MLProgramAttributes.h - Attribute Classes ----------------*- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_MLPROGRAM_IR_MLPROGRAMATTRIBUTES_H_
 #define MLIR_DIALECT_MLPROGRAM_IR_MLPROGRAMATTRIBUTES_H_
 
 #include "mlir/IR/Attributes.h"
+#include "mlir/IR/BuiltinAttributeInterfaces.h"
 
 //===----------------------------------------------------------------------===//
 // Tablegen Attribute Declarations
 //===----------------------------------------------------------------------===//
 
 #define GET_ATTRDEF_CLASSES
 #include "mlir/Dialect/MLProgram/IR/MLProgramAttributes.h.inc"
 
 #endif // MLIR_DIALECT_MLPROGRAM_IR_MLPROGRAMATTRIBUTES_H_

mlir/include/mlir/Dialect/MLProgram/IR/MLProgramAttributes.td

 //===- MLProgramAttributed.td - Attr definitions -----------*- tablegen -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLPROGRAM_ATTRIBUTES
 #define MLPROGRAM_ATTRIBUTES
 
 include "mlir/IR/AttrTypeBase.td"
+include "mlir/IR/BuiltinAttributeInterfaces.td"
 include "mlir/Dialect/MLProgram/IR/MLProgramBase.td"
 
 // Base class for MLProgram dialect attributes.
 class MLProgram_Attr<string name, list<Trait> traits = []>
     : AttrDef<MLProgram_Dialect, name, traits> {
   let mnemonic = ?;
 }
 
 //===----------------------------------------------------------------------===//
 // ExternAttr
 //===----------------------------------------------------------------------===//
 
-def MLProgram_ExternAttr : MLProgram_Attr<"Extern"> {
+def MLProgram_ExternAttr : MLProgram_Attr<"Extern", [TypedAttrInterface]> {
   let summary = "Value used for a global signalling external resolution";
   let description = [{
   When used as the value for a GlobalOp, this indicates that the actual
   value should be resolved externally in an implementation defined manner.
   The `sym_name` of the global is the key for locating the value.
 
   Examples:
 

mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td

 // from the SPIR-V binary format in order to utilize MLIR mechanisms.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_SPIRV_IR_STRUCTURE_OPS
 #define MLIR_DIALECT_SPIRV_IR_STRUCTURE_OPS
 
 include "mlir/Dialect/SPIRV/IR/SPIRVBase.td"
+include "mlir/IR/BuiltinAttributeInterfaces.td"
 include "mlir/IR/FunctionInterfaces.td"
 include "mlir/IR/OpAsmInterface.td"
 include "mlir/IR/SymbolInterfaces.td"
 include "mlir/Interfaces/CallInterfaces.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 
 // -----
 
let description = [{
     ```mlir
     spv.SpecConstant @spec_const1 = true
     spv.SpecConstant @spec_const2 spec_id(5) = 42 : i32
     ```
   }];
 
   let arguments = (ins
     StrAttr:$sym_name,
-    AnyAttr:$default_value
+    TypedAttrInterface:$default_value
   );
 
   let results = (outs);
 
   let hasOpcode = 0;
 
   let autogenSerialization = 0;
 }

mlir/include/mlir/IR/AttrTypeBase.td

class AttrDef<Dialect dialect, string name, list<Trait> traits = [],
   //
   // Format: $_self will be expanded to the attribute.
   //
   // For example, `$_self.getValue().getSExtValue()` for `IntegerAttr val` will
   // expand to `getAttrOfType<IntegerAttr>("val").getValue().getSExtValue()`.
   let convertFromStorage = "$_self.cast<" # dialect.cppNamespace #
                                  "::" # cppClassName # ">()";
 
-  // A code block used to build the value 'Type' of an Attribute when
-  // initializing its storage instance. This field is optional, and if not
-  // present the attribute will have its value type set to `NoneType`. This code
-  // block may reference any of the attributes parameters via
-  // `$_<parameter-name`. If one of the parameters of the attribute is of type
-  // `AttributeSelfTypeParameter`, this field is ignored.
-  code typeBuilder = ?;
-
   // The predicate for when this def is used as a constraint.
   let predicate = CPred<"$_self.isa<" # dialect.cppNamespace #
                                  "::" # cppClassName # ">()">;
 }
 
 // Define a new type, named `name`, belonging to `dialect` that inherits from
 // the given C++ base class.
 class TypeDef<Dialect dialect, string name, list<Trait> traits = [],
class AttrOrTypeParameter<string type, string desc, string accessorType = ""> {
   bit isOptional = 0;
   // Provide a default value for the parameter. Parameters with default values
   // are considered optional. If a value was not parsed for the parameter, it
   // will be set to the default value. Parameters equal to their default values
   // are elided when printing. Equality is checked using the `comparator` field,
   // which by default is the C++ equality operator. The current MLIR context is
   // made available through `$_ctxt`, e.g., for constructing default values for
   // attributes and types.
-  string defaultValue = ?;
+  string defaultValue = "";
 }
 class AttrParameter<string type, string desc, string accessorType = "">
  : AttrOrTypeParameter<type, desc, accessorType>;
 class TypeParameter<string type, string desc, string accessorType = "">
  : AttrOrTypeParameter<type, desc, accessorType>;
 
 // For StringRefs, which require allocation.
 class StringRefParameter<string desc = ""> :
class DefaultValuedParameter<string type, string value, string desc = ""> :
   let isOptional = 1;
   let defaultValue = value;
 }
 
 // This is a special parameter used for AttrDefs that represents a `mlir::Type`
 // that is also used as the value `Type` of the attribute. Only one parameter
 // of the attribute may be of this type.
 class AttributeSelfTypeParameter<string desc,
-                                 string derivedType = "::mlir::Type"> :
-    AttrOrTypeParameter<derivedType, desc> {}
+                                 string derivedType = "::mlir::Type",
+                                 string typeBuilder = ""> :
+    AttrOrTypeParameter<derivedType, desc> {
+  let defaultValue = !if(!and(!empty(typeBuilder),
+                              !eq(derivedType, "::mlir::Type")),
+                         "::mlir::NoneType::get($_ctxt)", typeBuilder);
+}
 
 #endif // ATTRTYPEBASE_TD

mlir/include/mlir/IR/AttributeSupport.h

 
 /// Base storage class appearing in an attribute. Derived storage classes should
 /// only be constructed within the context of the AttributeUniquer.
 class alignas(8) AttributeStorage : public StorageUniquer::BaseStorage {
   friend detail::AttributeUniquer;
   friend StorageUniquer;
 
 public:
-  /// Get the type of this attribute.
-  Type getType() const { return type; }
-
   /// Return the abstract descriptor for this attribute.
   const AbstractAttribute &getAbstractAttribute() const {
     assert(abstractAttribute && "Malformed attribute storage object.");
     return *abstractAttribute;
   }
 
 protected:
-  /// Construct a new attribute storage instance with the given type.
-  /// Note: All attributes require a valid type. If no type is provided here,
-  ///       the type of the attribute will automatically default to NoneType
-  ///       upon initialization in the uniquer.
-  AttributeStorage(Type type = nullptr) : type(type) {}
-
-  /// Set the type of this attribute.
-  void setType(Type newType) { type = newType; }
-
   /// Set the abstract attribute for this storage instance. This is used by the
   /// AttributeUniquer when initializing a newly constructed storage object.
   void initializeAbstractAttribute(const AbstractAttribute &abstractAttr) {
     abstractAttribute = &abstractAttr;
   }
 
   /// Default initialization for attribute storage classes that require no
   /// additional initialization.
   void initialize(MLIRContext *context) {}
 
 private:
-  /// The type of the attribute value.
-  Type type;
-
   /// The abstract descriptor for this attribute.
   const AbstractAttribute *abstractAttribute = nullptr;
 };
 
 /// Default storage type for attributes that require no additional
 /// initialization or storage.
 using DefaultAttributeStorage = AttributeStorage;
 

mlir/include/mlir/IR/Attributes.h

public:
 
   // Support dyn_cast'ing Attribute to itself.
   static bool classof(Attribute) { return true; }
 
   /// Return a unique identifier for the concrete attribute type. This is used
   /// to support dynamic type casting.
   TypeID getTypeID() { return impl->getAbstractAttribute().getTypeID(); }
 
-  /// Return the type of this attribute.
-  Type getType() const { return impl->getType(); }
-
   /// Return the context this attribute belongs to.
   MLIRContext *getContext() const;
 
   /// Get the dialect this attribute is registered to.
   Dialect &getDialect() const {
     return impl->getAbstractAttribute().getDialect();
   }
 

mlir/include/mlir/IR/BuiltinAttributeInterfaces.h

 //===- BuiltinAttributeInterfaces.h - Builtin Attr Interfaces ---*- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_IR_BUILTINATTRIBUTEINTERFACES_H
 #define MLIR_IR_BUILTINATTRIBUTEINTERFACES_H
 
 #include "mlir/IR/AffineMap.h"
 #include "mlir/IR/Attributes.h"
+#include "mlir/IR/BuiltinTypeInterfaces.h"
 #include "mlir/IR/Types.h"
 #include "mlir/Support/LogicalResult.h"
 #include "llvm/ADT/Any.h"
 #include "llvm/Support/raw_ostream.h"
 #include <complex>
 
 namespace mlir {
-class ShapedType;
 
 //===----------------------------------------------------------------------===//
 // ElementsAttr
 //===----------------------------------------------------------------------===//
 namespace detail {
 /// This class provides support for indexing into the element range of an
 /// ElementsAttr. It is used to opaquely wrap either a contiguous range, via
 /// `ElementsAttrIndexer::contiguous`, or a non-contiguous range, via

 };
 
 /// This class provides iterator utilities for an ElementsAttr range.
 template <typename IteratorT>
 class ElementsAttrRange : public llvm::iterator_range<IteratorT> {
 public:
   using reference = typename IteratorT::reference;
 
-  ElementsAttrRange(Type shapeType,
+  ElementsAttrRange(ShapedType shapeType,
                     const llvm::iterator_range<IteratorT> &range)
       : llvm::iterator_range<IteratorT>(range), shapeType(shapeType) {}
-  ElementsAttrRange(Type shapeType, IteratorT beginIt, IteratorT endIt)
+  ElementsAttrRange(ShapedType shapeType, IteratorT beginIt, IteratorT endIt)
       : ElementsAttrRange(shapeType, llvm::make_range(beginIt, endIt)) {}
 
   /// Return the value at the given index.
   reference operator[](ArrayRef<uint64_t> index) const;
   reference operator[](uint64_t index) const {
     return *std::next(this->begin(), index);
   }
 
   /// Return the size of this range.
   size_t size() const { return llvm::size(*this); }
 
 private:
   /// The shaped type of the parent ElementsAttr.
-  Type shapeType;
+  ShapedType shapeType;
 };
 
 } // namespace detail
 
 //===----------------------------------------------------------------------===//
 // MemRefLayoutAttrInterface
 //===----------------------------------------------------------------------===//
 

mlir/include/mlir/IR/BuiltinAttributeInterfaces.td

(ins "::mlir::TypeID":$elementID), [{}], /*defaultImplementation=*/[{
         elementID, /*isContiguous=*/std::false_type());
     }]>,
     InterfaceMethod<[{
       Returns true if the attribute elements correspond to a splat, i.e. that
       all elements of the attribute are the same value.
     }], "bool", "isSplat", (ins), /*defaultImplementation=*/[{}], [{
         // By default, only check for a single element splat.
         return $_attr.getNumElements() == 1;
-    }]>
+    }]>,
+    InterfaceMethod<[{
+      Returns the shaped type of the elements attribute.
+    }], "::mlir::ShapedType", "getType", (ins)>
   ];
 
   string ElementsAttrInterfaceAccessors = [{
     /// Return the number of elements held by this attribute.
     int64_t size() const { return getNumElements(); }
 
     /// Return if the attribute holds no elements.
     bool empty() const { return size() == 0; }
auto value_begin() const {
       return $_attr.value_begin_impl(OverloadToken<T>());
     }
 
     /// Return the elements of this attribute as a value of type 'T'.
     template <typename T>
     auto getValues() const {
       auto beginIt = $_attr.template value_begin<T>();
       return detail::ElementsAttrRange<decltype(beginIt)>(
-        Attribute($_attr).getType(), beginIt, std::next(beginIt, size()));
+        $_attr.getType(), beginIt, std::next(beginIt, size()));
     }
   }] # ElementsAttrInterfaceAccessors;
 
   let extraClassDeclaration = [{
     template <typename T>
     using iterator = detail::ElementsAttrIterator<T>;
     template <typename T>
     using iterator_range = detail::ElementsAttrRange<iterator<T>>;
 
     //===------------------------------------------------------------------===//
     // Accessors
     //===------------------------------------------------------------------===//
 
-    /// Return the type of this attribute.
-    ShapedType getType() const;
-
     /// Return the element type of this ElementsAttr.
     Type getElementType() const { return getElementType(*this); }
-    static Type getElementType(Attribute elementsAttr);
+    static Type getElementType(ElementsAttr elementsAttr);
 
     /// Return if the given 'index' refers to a valid element in this attribute.
     bool isValidIndex(ArrayRef<uint64_t> index) const {
       return isValidIndex(*this, index);
     }
     static bool isValidIndex(ShapedType type, ArrayRef<uint64_t> index);
-    static bool isValidIndex(Attribute elementsAttr, ArrayRef<uint64_t> index);
+    static bool isValidIndex(ElementsAttr elementsAttr,
+                             ArrayRef<uint64_t> index);
 
     /// Return the 1 dimensional flattened row-major index from the given
     /// multi-dimensional index.
     uint64_t getFlattenedIndex(ArrayRef<uint64_t> index) const {
       return getFlattenedIndex(*this, index);
     }
     static uint64_t getFlattenedIndex(Type type,
                                       ArrayRef<uint64_t> index);
-    static uint64_t getFlattenedIndex(Attribute elementsAttr,
+    static uint64_t getFlattenedIndex(ElementsAttr elementsAttr,
                                       ArrayRef<uint64_t> index) {
       return getFlattenedIndex(elementsAttr.getType(), index);
     }
 
     /// Returns the number of elements held by this attribute.
     int64_t getNumElements() const { return getNumElements(*this); }
-    static int64_t getNumElements(Attribute elementsAttr);
+    static int64_t getNumElements(ElementsAttr elementsAttr);
 
     //===------------------------------------------------------------------===//
     // Value Iteration
     //===------------------------------------------------------------------===//
 
     template <typename T>
     using DerivedAttrValueCheckT =
         typename std::enable_if_t<std::is_base_of<Attribute, T>::value &&
let extraClassDeclaration = [{
     T getSplatValue() const {
       assert(isSplat() && "expected splat attribute");
       return *value_begin<T>();
     }
 
     /// Return the elements of this attribute as a value of type 'T'.
     template <typename T>
     DefaultValueCheckT<T, iterator_range<T>> getValues() const {
-      return {Attribute::getType(), value_begin<T>(), value_end<T>()};
+      return {getType(), value_begin<T>(), value_end<T>()};
     }
     template <typename T>
     DefaultValueCheckT<T, iterator<T>> value_begin() const;
     template <typename T>
     DefaultValueCheckT<T, iterator<T>> value_end() const {
       return iterator<T>({}, size());
     }
 
     /// Return the held element values a range of T, where T is a derived
     /// attribute type.
     template <typename T>
     using DerivedAttrValueIterator =
       llvm::mapped_iterator<iterator<Attribute>, T (*)(Attribute)>;
     template <typename T>
     using DerivedAttrValueIteratorRange =
       detail::ElementsAttrRange<DerivedAttrValueIterator<T>>;
     template <typename T, typename = DerivedAttrValueCheckT<T>>
     DerivedAttrValueIteratorRange<T> getValues() const {
       auto castFn = [](Attribute attr) { return attr.template cast<T>(); };
-      return {Attribute::getType(), llvm::map_range(getValues<Attribute>(),
-                             static_cast<T (*)(Attribute)>(castFn))};
+      return {getType(), llvm::map_range(getValues<Attribute>(),
     }
     template <typename T, typename = DerivedAttrValueCheckT<T>>
     DerivedAttrValueIterator<T> value_begin() const {
       return getValues<T>().begin();
     }
     template <typename T, typename = DerivedAttrValueCheckT<T>>
     DerivedAttrValueIterator<T> value_end() const {
       return {value_end<Attribute>(), nullptr};
     }
 
     //===------------------------------------------------------------------===//
     // Failable Value Iteration
 
     /// If this attribute supports iterating over element values of type `T`,
     /// return the iterable range. Otherwise, return llvm::None.
     template <typename T>
     DefaultValueCheckT<T, Optional<iterator_range<T>>> tryGetValues() const {
-      if (Optional<iterator<T>> beginIt = try_value_begin<T>()) {
-        return iterator_range<T>(Attribute::getType(), *beginIt,
-                                 value_end<T>());
-      }
+      if (Optional<iterator<T>> beginIt = try_value_begin<T>())
+        return iterator_range<T>(getType(), *beginIt, value_end<T>());
       return llvm::None;
     }
     template <typename T>
     DefaultValueCheckT<T, Optional<iterator<T>>> try_value_begin() const;
 
     /// If this attribute supports iterating over element values of type `T`,
     /// return the iterable range. Otherwise, return llvm::None.
     template <typename T, typename = DerivedAttrValueCheckT<T>>
     Optional<DerivedAttrValueIteratorRange<T>> tryGetValues() const {
       auto values = tryGetValues<Attribute>();
       if (!values)
         return llvm::None;
 
       auto castFn = [](Attribute attr) { return attr.template cast<T>(); };
       return DerivedAttrValueIteratorRange<T>(
-        Attribute::getType(),
+        getType(),
         llvm::map_range(*values, static_cast<T (*)(Attribute)>(castFn))
       );
     }
     template <typename T, typename = DerivedAttrValueCheckT<T>>
     Optional<DerivedAttrValueIterator<T>> try_value_begin() const {
       if (auto values = tryGetValues<T>())
         return values->begin();
       return llvm::None;
InterfaceMethod<
       [{
         return ::mlir::detail::verifyAffineMapAsLayout($_attr.getAffineMap(),
                                                        shape, emitError);
       }]
     >
   ];
 }
 
+//===----------------------------------------------------------------------===//
+// TypedAttrInterface
+//===----------------------------------------------------------------------===//
+
+def TypedAttrInterface : AttrInterface<"TypedAttr"> {
+  let cppNamespace = "::mlir";
+
+  let description = [{
+    This interface is used for attributes with types. The type of an attribute
+    is understood to represent the type of the data contained in the attribute
+    and is often used to represent the type of a value with this data.
+  }];
+
+  let methods = [InterfaceMethod<
+    "Get the attribute's type", "::mlir::Type", "getType", (ins)>];
+}
+
 #endif // MLIR_IR_BUILTINATTRIBUTEINTERFACES_TD_

mlir/include/mlir/IR/BuiltinAttributes.h

 class BoolAttr;
 class DenseIntElementsAttr;
 class FlatSymbolRefAttr;
 class FunctionType;
 class IntegerSet;
 class IntegerType;
 class Location;
 class Operation;
-class ShapedType;
 
 //===----------------------------------------------------------------------===//
 // Elements Attributes
 //===----------------------------------------------------------------------===//
 
 namespace detail {
 /// Pair of raw pointer and a boolean flag of whether the pointer holds a splat,
 using DenseIterPtrAndSplat = std::pair<const char *, bool>;
using IntFloatValueTemplateCheckT =
                                std::numeric_limits<T>::is_integer) ||
                               is_valid_cpp_fp_type<T>::value>::type;
   template <typename T, typename = IntFloatValueTemplateCheckT<T>>
   iterator_range_impl<ElementIterator<T>> getValues() const {
     assert(isValidIntOrFloat(sizeof(T), std::numeric_limits<T>::is_integer,
                              std::numeric_limits<T>::is_signed));
     const char *rawData = getRawData().data();
     bool splat = isSplat();
-    return {Attribute::getType(), ElementIterator<T>(rawData, splat, 0),
+    return {getType(), ElementIterator<T>(rawData, splat, 0),
             ElementIterator<T>(rawData, splat, getNumElements())};
   }
   template <typename T, typename = IntFloatValueTemplateCheckT<T>>
   ElementIterator<T> value_begin() const {
     assert(isValidIntOrFloat(sizeof(T), std::numeric_limits<T>::is_integer,
                              std::numeric_limits<T>::is_signed));
     return ElementIterator<T>(getRawData().data(), isSplat(), 0);
   }
using ComplexValueTemplateCheckT =
                                is_valid_cpp_fp_type<ElementT>::value)>::type;
   template <typename T, typename ElementT = typename T::value_type,
             typename = ComplexValueTemplateCheckT<T, ElementT>>
   iterator_range_impl<ElementIterator<T>> getValues() const {
     assert(isValidComplex(sizeof(T), std::numeric_limits<ElementT>::is_integer,
                           std::numeric_limits<ElementT>::is_signed));
     const char *rawData = getRawData().data();
     bool splat = isSplat();
-    return {Attribute::getType(), ElementIterator<T>(rawData, splat, 0),
+    return {getType(), ElementIterator<T>(rawData, splat, 0),
             ElementIterator<T>(rawData, splat, getNumElements())};
   }
   template <typename T, typename ElementT = typename T::value_type,
             typename = ComplexValueTemplateCheckT<T, ElementT>>
   ElementIterator<T> value_begin() const {
     assert(isValidComplex(sizeof(T), std::numeric_limits<ElementT>::is_integer,
                           std::numeric_limits<ElementT>::is_signed));
     return ElementIterator<T>(getRawData().data(), isSplat(), 0);
public:
   template <typename T>
   using StringRefValueTemplateCheckT =
       typename std::enable_if<std::is_same<T, StringRef>::value>::type;
   template <typename T, typename = StringRefValueTemplateCheckT<T>>
   iterator_range_impl<ElementIterator<StringRef>> getValues() const {
     auto stringRefs = getRawStringData();
     const char *ptr = reinterpret_cast<const char *>(stringRefs.data());
     bool splat = isSplat();
-    return {Attribute::getType(), ElementIterator<StringRef>(ptr, splat, 0),
+    return {getType(), ElementIterator<StringRef>(ptr, splat, 0),
             ElementIterator<StringRef>(ptr, splat, getNumElements())};
   }
   template <typename T, typename = StringRefValueTemplateCheckT<T>>
   ElementIterator<StringRef> value_begin() const {
     const char *ptr = reinterpret_cast<const char *>(getRawStringData().data());
     return ElementIterator<StringRef>(ptr, isSplat(), 0);
   }
   template <typename T, typename = StringRefValueTemplateCheckT<T>>
   ElementIterator<StringRef> value_end() const {
     const char *ptr = reinterpret_cast<const char *>(getRawStringData().data());
     return ElementIterator<StringRef>(ptr, isSplat(), getNumElements());
   }
 
   /// Return the held element values as a range of Attributes.
   template <typename T>
   using AttributeValueTemplateCheckT =
       typename std::enable_if<std::is_same<T, Attribute>::value>::type;
   template <typename T, typename = AttributeValueTemplateCheckT<T>>
   iterator_range_impl<AttributeElementIterator> getValues() const {
-    return {Attribute::getType(), value_begin<Attribute>(),
-            value_end<Attribute>()};
+    return {getType(), value_begin<Attribute>(), value_end<Attribute>()};
   }
   template <typename T, typename = AttributeValueTemplateCheckT<T>>
   AttributeElementIterator value_begin() const {
     return AttributeElementIterator(*this, 0);
   }
   template <typename T, typename = AttributeValueTemplateCheckT<T>>
   AttributeElementIterator value_end() const {
     return AttributeElementIterator(*this, getNumElements());
using llvm::mapped_iterator_base<DerivedAttributeElementIterator<T>,
                                      T>::mapped_iterator_base;
 
     /// Map the element to the iterator result type.
     T mapElement(Attribute attr) const { return attr.cast<T>(); }
   };
   template <typename T, typename = DerivedAttrValueTemplateCheckT<T>>
   iterator_range_impl<DerivedAttributeElementIterator<T>> getValues() const {
     using DerivedIterT = DerivedAttributeElementIterator<T>;
-    return {Attribute::getType(), DerivedIterT(value_begin<Attribute>()),
+    return {getType(), DerivedIterT(value_begin<Attribute>()),
             DerivedIterT(value_end<Attribute>())};
   }
   template <typename T, typename = DerivedAttrValueTemplateCheckT<T>>
   DerivedAttributeElementIterator<T> value_begin() const {
     return {value_begin<Attribute>()};
   }
   template <typename T, typename = DerivedAttrValueTemplateCheckT<T>>
   DerivedAttributeElementIterator<T> value_end() const {
     return {value_end<Attribute>()};
   }
 
   /// Return the held element values as a range of bool. The element type of
   /// this attribute must be of integer type of bitwidth 1.
   template <typename T>
   using BoolValueTemplateCheckT =
       typename std::enable_if<std::is_same<T, bool>::value>::type;
   template <typename T, typename = BoolValueTemplateCheckT<T>>
   iterator_range_impl<BoolElementIterator> getValues() const {
     assert(isValidBool() && "bool is not the value of this elements attribute");
-    return {Attribute::getType(), BoolElementIterator(*this, 0),
+    return {getType(), BoolElementIterator(*this, 0),
             BoolElementIterator(*this, getNumElements())};
   }
   template <typename T, typename = BoolValueTemplateCheckT<T>>
   BoolElementIterator value_begin() const {
     assert(isValidBool() && "bool is not the value of this elements attribute");
     return BoolElementIterator(*this, 0);
   }
   template <typename T, typename = BoolValueTemplateCheckT<T>>
   BoolElementIterator value_end() const {
     assert(isValidBool() && "bool is not the value of this elements attribute");
     return BoolElementIterator(*this, getNumElements());
   }
 
   /// Return the held element values as a range of APInts. The element type of
   /// this attribute must be of integer type.
   template <typename T>
   using APIntValueTemplateCheckT =
       typename std::enable_if<std::is_same<T, APInt>::value>::type;
   template <typename T, typename = APIntValueTemplateCheckT<T>>
   iterator_range_impl<IntElementIterator> getValues() const {
     assert(getElementType().isIntOrIndex() && "expected integral type");
-    return {Attribute::getType(), raw_int_begin(), raw_int_end()};
+    return {getType(), raw_int_begin(), raw_int_end()};
   }
   template <typename T, typename = APIntValueTemplateCheckT<T>>
   IntElementIterator value_begin() const {
     assert(getElementType().isIntOrIndex() && "expected integral type");
     return raw_int_begin();
   }
   template <typename T, typename = APIntValueTemplateCheckT<T>>
   IntElementIterator value_end() const {
inline bool operator==(StringAttr lhs, StringRef rhs) {
   return lhs.getValue() == rhs;
 }
 inline bool operator!=(StringAttr lhs, StringRef rhs) { return !(lhs == rhs); }
 inline bool operator==(StringRef lhs, StringAttr rhs) {
   return rhs.getValue() == lhs;
 }
 inline bool operator!=(StringRef lhs, StringAttr rhs) { return !(lhs == rhs); }
 
-inline Type StringAttr::getType() const { return Attribute::getType(); }
-
 } // namespace mlir
 
 //===----------------------------------------------------------------------===//
 // Attribute Utilities
 //===----------------------------------------------------------------------===//
 
 namespace llvm {
 

mlir/include/mlir/IR/BuiltinAttributes.td

AttrBuilderWithInferredContext<(ins "AffineMap":$value), [{
       return $_get(value.getContext(), value);
     }]>
   ];
   let extraClassDeclaration = [{
     using ValueType = AffineMap;
     AffineMap getAffineMap() const { return getValue(); }
   }];
   let skipDefaultBuilders = 1;
-  let typeBuilder = "IndexType::get($_value.getContext())";
 }
 
 //===----------------------------------------------------------------------===//
 // ArrayAttr
 //===----------------------------------------------------------------------===//
 
 def Builtin_ArrayAttr : Builtin_Attr<"Array", [
     DeclareAttrInterfaceMethods<SubElementAttrInterface,
auto getAsValueRange() const {
       return llvm::map_range(getAsRange<AttrTy>(), [](AttrTy attr) {
         return static_cast<UnderlyingTy>(attr.getValue());
       });
     }
   }];
 }
 
 //===----------------------------------------------------------------------===//
-// DenseIntOrFPElementsAttr
+// DenseArrayBaseAttr
 //===----------------------------------------------------------------------===//
 
 def Builtin_DenseArrayBase : Builtin_Attr<
-    "DenseArrayBase", [ElementsAttrInterface]> {
+    "DenseArrayBase", [ElementsAttrInterface, TypedAttrInterface]> {
   let summary = "A dense array of i8, i16, i32, i64, f32, or f64.";
   let description = [{
     A dense array attribute is an attribute that represents a dense array of
     primitive element types. Contrary to DenseIntOrFPElementsAttr this is a
     flat unidimensional array which does not have a storage optimization for
     splat. This allows to expose the raw array through a C++ API as
     `ArrayRef<T>`. This is the base class attribute, the actual access is
     intended to be managed through the subclasses `DenseI8ArrayAttr`,
using ContiguousIterableTypesT =
         std::tuple<int8_t, int16_t, int32_t, int64_t, float, double>;
     const int8_t *value_begin_impl(OverloadToken<int8_t>) const;
     const int16_t *value_begin_impl(OverloadToken<int16_t>) const;
     const int32_t *value_begin_impl(OverloadToken<int32_t>) const;
     const int64_t *value_begin_impl(OverloadToken<int64_t>) const;
     const float *value_begin_impl(OverloadToken<float>) const;
     const double *value_begin_impl(OverloadToken<double>) const;
 
-    /// Methods to support type inquiry through isa, cast, and dyn_cast.
+    /// Returns the shaped type.
+    ShapedType getType() const;
+    /// Returns the element type.
     EltType getElementType() const;
+
     /// Printer for the short form: will dispatch to the appropriate subclass.
     void print(AsmPrinter &printer) const;
     void print(raw_ostream &os) const;
     /// Print the short form `42, 100, -1` without any braces or prefix.
     void printWithoutBraces(raw_ostream &os) const;
   }];
   // Do not generate the storage class, we need to handle custom storage alignment.
   let genStorageClass = 0;
   let genAccessors = 0;
   let skipDefaultBuilders = 1;
 }
 
 //===----------------------------------------------------------------------===//
 // DenseIntOrFPElementsAttr
 //===----------------------------------------------------------------------===//
 
 def Builtin_DenseIntOrFPElementsAttr : Builtin_Attr<
-    "DenseIntOrFPElements", [ElementsAttrInterface], "DenseElementsAttr"
+    "DenseIntOrFPElements", [ElementsAttrInterface, TypedAttrInterface],
+    "DenseElementsAttr"
   > {
   let summary = "An Attribute containing a dense multi-dimensional array of "
                 "integer or floating-point values";
   let description = [{
     Syntax:
 
     ```
     tensor-literal ::= integer-literal | float-literal | bool-literal | [] | [tensor-literal (, tensor-literal)* ]
def Builtin_DenseIntOrFPElementsAttr : Builtin_Attr<
   let skipDefaultBuilders = 1;
 }
 
 //===----------------------------------------------------------------------===//
 // DenseStringElementsAttr
 //===----------------------------------------------------------------------===//
 
 def Builtin_DenseStringElementsAttr : Builtin_Attr<
-    "DenseStringElements", [ElementsAttrInterface], "DenseElementsAttr"
+    "DenseStringElements", [ElementsAttrInterface, TypedAttrInterface],
+    "DenseElementsAttr"
   > {
   let summary = "An Attribute containing a dense multi-dimensional array of "
                 "strings";
   let description = [{
     Syntax:
 
     ```
     dense-string-elements-attribute ::= `dense` `<` attribute-value `>` `:`
def Builtin_DictionaryAttr : Builtin_Attr<"Dictionary", [
   }];
   let skipDefaultBuilders = 1;
 }
 
 //===----------------------------------------------------------------------===//
 // FloatAttr
 //===----------------------------------------------------------------------===//
 
-def Builtin_FloatAttr : Builtin_Attr<"Float"> {
+def Builtin_FloatAttr : Builtin_Attr<"Float", [TypedAttrInterface]> {
   let summary = "An Attribute containing a floating-point value";
   let description = [{
     Syntax:
 
     ```
     float-attribute ::= (float-literal (`:` float-type)?)
                       | (hexadecimal-literal `:` float-type)
     ```
def Builtin_FloatAttr : Builtin_Attr<"Float", [TypedAttrInterface]> {
   let genVerifyDecl = 1;
   let skipDefaultBuilders = 1;
 }
 
 //===----------------------------------------------------------------------===//
 // IntegerAttr
 //===----------------------------------------------------------------------===//
 
-def Builtin_IntegerAttr : Builtin_Attr<"Integer"> {
+def Builtin_IntegerAttr : Builtin_Attr<"Integer", [TypedAttrInterface]> {
   let summary = "An Attribute containing a integer value";
   let description = [{
     Syntax:
 
     ```
     integer-attribute ::= (integer-literal ( `:` (index-type | integer-type) )?)
                           | `true` | `false`
     ```
def Builtin_IntegerSetAttr : Builtin_Attr<"IntegerSet"> {
   let extraClassDeclaration = "using ValueType = IntegerSet;";
   let skipDefaultBuilders = 1;
 }
 
 //===----------------------------------------------------------------------===//
 // OpaqueAttr
 //===----------------------------------------------------------------------===//
 
-def Builtin_OpaqueAttr : Builtin_Attr<"Opaque"> {
+def Builtin_OpaqueAttr : Builtin_Attr<"Opaque", [TypedAttrInterface]> {
   let summary = "An opaque representation of another Attribute";
   let description = [{
     Syntax:
 
     ```
     opaque-attribute ::= dialect-namespace `<` attr-data `>`
     ```
 
def Builtin_OpaqueAttr : Builtin_Attr<"Opaque", [TypedAttrInterface]> {
   let skipDefaultBuilders = 1;
 }
 
 //===----------------------------------------------------------------------===//
 // OpaqueElementsAttr
 //===----------------------------------------------------------------------===//
 
 def Builtin_OpaqueElementsAttr : Builtin_Attr<
-    "OpaqueElements", [ElementsAttrInterface]
+    "OpaqueElements", [ElementsAttrInterface, TypedAttrInterface]
   > {
   let summary = "An opaque representation of a multi-dimensional array";
   let description = [{
     Syntax:
 
     ```
     opaque-elements-attribute ::= `opaque` `<` dialect-namespace  `,`
                                   hex-string-literal `>` `:`
def Builtin_OpaqueElementsAttr : Builtin_Attr<
   let skipDefaultBuilders = 1;
 }
 
 //===----------------------------------------------------------------------===//
 // SparseElementsAttr
 //===----------------------------------------------------------------------===//
 
 def Builtin_SparseElementsAttr : Builtin_Attr<
-    "SparseElements", [ElementsAttrInterface]
+    "SparseElements", [ElementsAttrInterface, TypedAttrInterface]
   > {
   let summary = "An opaque representation of a multi-dimensional array";
   let description = [{
     Syntax:
 
     ```
     sparse-elements-attribute ::= `sparse` `<` attribute-value `,`
                                   attribute-value `>` `:`
def Builtin_SparseElementsAttr : Builtin_Attr<
   let genVerifyDecl = 1;
   let skipDefaultBuilders = 1;
 }
 
 //===----------------------------------------------------------------------===//
 // StringAttr
 //===----------------------------------------------------------------------===//
 
-def Builtin_StringAttr : Builtin_Attr<"String"> {
+def Builtin_StringAttr : Builtin_Attr<"String", [TypedAttrInterface]> {
   let summary = "An Attribute containing a string";
   let description = [{
     Syntax:
 
     ```
     string-attribute ::= string-literal (`:` type)?
     ```
 

mlir/include/mlir/IR/BuiltinTypeInterfaces.h

+//===- BuiltinTypeInterfaces.h - Builtin Type Interfaces --------*- C++ -*-===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_IR_BUILTINTYPEINTERFACES_H
+#define MLIR_IR_BUILTINTYPEINTERFACES_H
+
+#include "mlir/IR/BuiltinTypeInterfaces.h.inc"
+
+#endif // MLIR_IR_BUILTINTYPEINTERFACES_H

mlir/include/mlir/IR/BuiltinTypes.h

 //===- BuiltinTypes.h - MLIR Builtin Type Classes ---------------*- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_IR_BUILTINTYPES_H
 #define MLIR_IR_BUILTINTYPES_H
 
-#include "BuiltinAttributeInterfaces.h"
-#include "SubElementInterfaces.h"
+#include "mlir/IR/BuiltinAttributeInterfaces.h"
+#include "mlir/IR/BuiltinTypeInterfaces.h"
+#include "mlir/IR/SubElementInterfaces.h"
 
 namespace llvm {
 class BitVector;
 struct fltSemantics;
 } // namespace llvm
 
 //===----------------------------------------------------------------------===//
 // Tablegen Interface Declarations
 //===----------------------------------------------------------------------===//
 
-#include "mlir/IR/BuiltinTypeInterfaces.h.inc"
-
 namespace mlir {
 class AffineExpr;
 class AffineMap;
 class FloatType;
 class IndexType;
 class IntegerType;
 class StringAttr;
 class TypeRange;

mlir/lib/CAPI/IR/IR.cpp

MlirAttribute mlirAttributeParseGet(MlirContext context, MlirStringRef attr) {
   return wrap(mlir::parseAttribute(unwrap(attr), unwrap(context)));
 }
 
 MlirContext mlirAttributeGetContext(MlirAttribute attribute) {
   return wrap(unwrap(attribute).getContext());
 }
 
 MlirType mlirAttributeGetType(MlirAttribute attribute) {
-  return wrap(unwrap(attribute).getType());
+  Attribute attr = unwrap(attribute);
+  if (auto typedAttr = attr.dyn_cast<TypedAttr>())
+    return wrap(typedAttr.getType());
+  return wrap(NoneType::get(attr.getContext()));
 }
 
 MlirTypeID mlirAttributeGetTypeID(MlirAttribute attr) {
   return wrap(unwrap(attr).getTypeID());
 }
 
 bool mlirAttributeEqual(MlirAttribute a1, MlirAttribute a2) {
   return unwrap(a1) == unwrap(a2);

mlir/lib/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.cpp

if (auto shapedType = srcType.dyn_cast<ShapedType>()) {
     if (shapedType.getNumElements() != 1)
       return failure();
     srcType = shapedType.getElementType();
   }
   if (!srcType.isIntOrIndexOrFloat())
     return failure();
 
   Attribute cstAttr = constOp.getValue();
-  if (cstAttr.getType().isa<ShapedType>())
-    cstAttr = cstAttr.cast<DenseElementsAttr>().getSplatValue<Attribute>();
+  if (auto elementsAttr = cstAttr.dyn_cast<DenseElementsAttr>())
+    cstAttr = elementsAttr.getSplatValue<Attribute>();
 
   Type dstType = getTypeConverter()->convertType(srcType);
   if (!dstType)
     return failure();
 
   // Floating-point types.
   if (srcType.isa<FloatType>()) {
     auto srcAttr = cstAttr.cast<FloatAttr>();

mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp

convertContractOpToMmaSync(vector::ContractionOp op,
   return success();
 }
 
 /// Convert a 2D splat ConstantOp to a SubgroupMmaConstantMatrix op.
 static void convertConstantOp(arith::ConstantOp op,
                               llvm::DenseMap<Value, Value> &valueMapping) {
   assert(constantSupportsMMAMatrixType(op));
   OpBuilder b(op);
-  Attribute splat =
-      op.getValue().cast<SplatElementsAttr>().getSplatValue<Attribute>();
+  auto splat =
+      op.getValue().cast<SplatElementsAttr>().getSplatValue<TypedAttr>();
   auto scalarConstant =
       b.create<arith::ConstantOp>(op.getLoc(), splat.getType(), splat);
   const char *fragType = inferFragType(op);
   auto vecType = op.getType().cast<VectorType>();
   gpu::MMAMatrixType type = gpu::MMAMatrixType::get(
       vecType.getShape(), vecType.getElementType(), llvm::StringRef(fragType));
   auto matrix = b.create<gpu::SubgroupMmaConstantMatrixOp>(op.getLoc(), type,
                                                            scalarConstant);

mlir/lib/Dialect/Arithmetic/IR/ArithmeticOps.cpp

if (!getValue().isa<IntegerAttr, FloatAttr, ElementsAttr>()) {
     return emitOpError(
         "value must be an integer, float, or elements attribute");
   }
   return success();
 }
 
 bool arith::ConstantOp::isBuildableWith(Attribute value, Type type) {
   // The value's type must be the same as the provided type.
-  if (value.getType() != type)
+  auto typedAttr = value.dyn_cast<TypedAttr>();
+  if (!typedAttr || typedAttr.getType() != type)
     return false;
   // Integer values must be signless.
   if (type.isa<IntegerType>() && !type.cast<IntegerType>().isSignless())
     return false;
   // Integer, float, and element attributes are buildable.
   return value.isa<IntegerAttr, FloatAttr, ElementsAttr>();
 }
 

mlir/lib/Dialect/Complex/IR/ComplexOps.cpp

 void ConstantOp::getAsmResultNames(
     function_ref<void(Value, StringRef)> setNameFn) {
   setNameFn(getResult(), "cst");
 }
 
 bool ConstantOp::isBuildableWith(Attribute value, Type type) {
   if (auto arrAttr = value.dyn_cast<ArrayAttr>()) {
     auto complexTy = type.dyn_cast<ComplexType>();
-    if (!complexTy)
+    if (!complexTy || arrAttr.size() != 2)
       return false;
     auto complexEltTy = complexTy.getElementType();
-    return arrAttr.size() == 2 && arrAttr[0].getType() == complexEltTy &&
-           arrAttr[1].getType() == complexEltTy;
+    auto re = arrAttr[0].dyn_cast<FloatAttr>();
+    auto im = arrAttr[1].dyn_cast<FloatAttr>();
+    return re && im && re.getType() == complexEltTy &&
+           im.getType() == complexEltTy;
   }
   return false;
 }
 
 LogicalResult ConstantOp::verify() {
   ArrayAttr arrayAttr = getValue();
   if (arrayAttr.size() != 2) {
     return emitOpError(
         "requires 'value' to be a complex constant, represented as array of "
         "two values");
   }
 
   auto complexEltTy = getType().getElementType();
-  if (complexEltTy != arrayAttr[0].getType() ||
-      complexEltTy != arrayAttr[1].getType()) {
+  auto re = arrayAttr[0].dyn_cast<FloatAttr>();
+  auto im = arrayAttr[1].dyn_cast<FloatAttr>();
+  if (!re || !im)
+    return emitOpError("requires attribute's elements to be float attributes");
+  if (complexEltTy != re.getType() || complexEltTy != im.getType()) {
     return emitOpError()
-           << "requires attribute's element types (" << arrayAttr[0].getType()
-           << ", " << arrayAttr[1].getType()
+           << "requires attribute's element types (" << re.getType() << ", "
+           << im.getType()
            << ") to match the element type of the op's return type ("
            << complexEltTy << ")";
   }
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // CreateOp

mlir/lib/Dialect/EmitC/IR/EmitC.cpp

 
 LogicalResult emitc::CallOp::verify() {
   // Callee must not be empty.
   if (getCallee().empty())
     return emitOpError("callee must not be empty");
 
   if (Optional<ArrayAttr> argsAttr = getArgs()) {
     for (Attribute arg : *argsAttr) {
-      if (arg.getType().isa<IndexType>()) {
-        int64_t index = arg.cast<IntegerAttr>().getInt();
+      auto intAttr = arg.dyn_cast<IntegerAttr>();
+      if (intAttr && intAttr.getType().isa<IndexType>()) {
+        int64_t index = intAttr.getInt();
         // Args with elements of type index must be in range
         // [0..operands.size).
         if ((index < 0) || (index >= static_cast<int64_t>(getNumOperands())))
           return emitOpError("index argument is out of range");
 
         // Args with elements of type ArrayAttr must have a type.
-      } else if (arg.isa<ArrayAttr>() && arg.getType().isa<NoneType>()) {
+      } else if (arg.isa<ArrayAttr>() /*&& arg.getType().isa<NoneType>()*/) {
+        // FIXME: Array attributes never have types
         return emitOpError("array argument has no type");
       }
     }
   }
 
   if (Optional<ArrayAttr> templateArgsAttr = getTemplateArgs()) {
     for (Attribute tArg : *templateArgsAttr) {
-      if (!tArg.isa<TypeAttr>() && !tArg.isa<IntegerAttr>() &&
-          !tArg.isa<FloatAttr>() && !tArg.isa<emitc::OpaqueAttr>())
+      if (!tArg.isa<TypeAttr, IntegerAttr, FloatAttr, emitc::OpaqueAttr>())
         return emitOpError("template argument has invalid type");
     }
   }
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // ConstantOp
 //===----------------------------------------------------------------------===//
 
 /// The constant op requires that the attribute's type matches the return type.
 LogicalResult emitc::ConstantOp::verify() {
-  Attribute value = getValueAttr();
+  TypedAttr value = getValueAttr();
   Type type = getType();
   if (!value.getType().isa<NoneType>() && type != value.getType())
     return emitOpError() << "requires attribute's type (" << value.getType()
                          << ") to match op's return type (" << type << ")";
   return success();
 }
 
 OpFoldResult emitc::ConstantOp::fold(ArrayRef<Attribute> operands) {

 }
 
 //===----------------------------------------------------------------------===//
 // VariableOp
 //===----------------------------------------------------------------------===//
 
 /// The variable op requires that the attribute's type matches the return type.
 LogicalResult emitc::VariableOp::verify() {
-  Attribute value = getValueAttr();
+  TypedAttr value = getValueAttr();
   Type type = getType();
   if (!value.getType().isa<NoneType>() && type != value.getType())
     return emitOpError() << "requires attribute's type (" << value.getType()
                          << ") to match op's return type (" << type << ")";
   return success();
 }
 
 //===----------------------------------------------------------------------===//
Attribute emitc::OpaqueAttr::parse(AsmParser &parser, Type type) {
   std::string value;
   SMLoc loc = parser.getCurrentLocation();
   if (parser.parseOptionalString(&value)) {
     parser.emitError(loc) << "expected string";
     return Attribute();
   }
   if (parser.parseGreater())
     return Attribute();
-  return get(parser.getContext(), value);
+
+  return get(parser.getContext(),
+             type ? type : NoneType::get(parser.getContext()), value);
 }
 
 void emitc::OpaqueAttr::print(AsmPrinter &printer) const {
   printer << "<\"";
   llvm::printEscapedString(getValue(), printer.getStream());
   printer << "\">";
 }
 

mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp

LogicalResult LLVM::ConstantOp::verify() {
   if (auto structType = getType().dyn_cast<LLVMStructType>()) {
     if (structType.getBody().size() != 2 ||
         structType.getBody()[0] != structType.getBody()[1]) {
       return emitError() << "expected struct type with two elements of the "
                             "same type, the type of a complex constant";
     }
 
     auto arrayAttr = getValue().dyn_cast<ArrayAttr>();
-    if (!arrayAttr || arrayAttr.size() != 2 ||
-        arrayAttr[0].getType() != arrayAttr[1].getType()) {
+    if (!arrayAttr || arrayAttr.size() != 2) {
       return emitOpError() << "expected array attribute with two elements, "
                               "representing a complex constant";
     }
+    auto re = arrayAttr[0].dyn_cast<TypedAttr>();
+    auto im = arrayAttr[1].dyn_cast<TypedAttr>();
+    if (!re || !im || re.getType() != im.getType()) {
+      return emitOpError()
+             << "expected array attribute with two elements of the same type";
+    }
 
     Type elementType = structType.getBody()[0];
     if (!elementType
              .isa<IntegerType, Float16Type, Float32Type, Float64Type>()) {
       return emitError()
              << "expected struct element types to be floating point type or "
                 "integer type";
     }

mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp

void yieldOutputs(ValueRange values) {
     Location loc = builder.getUnknownLoc();
     builder.create<YieldOp>(loc, values);
   }
 
   Value constant(const std::string &value) {
     OpBuilder builder = getBuilder();
     Location loc = builder.getUnknownLoc();
     Attribute valueAttr = parseAttribute(value, builder.getContext());
-    return builder.create<arith::ConstantOp>(loc, valueAttr.getType(),
-                                             valueAttr);
+    Type type = NoneType::get(builder.getContext());
+    if (auto typedAttr = valueAttr.dyn_cast<TypedAttr>())
+      type = typedAttr.getType();
+    return builder.create<arith::ConstantOp>(loc, type, valueAttr);
   }
 
   Value index(int64_t dim) {
     OpBuilder builder = getBuilder();
     return builder.create<IndexOp>(builder.getUnknownLoc(), dim);
   }
 
   Type getIntegerType(unsigned width) {

mlir/lib/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp

transform::PadOp::applyToOne(linalg::LinalgOp target,
   SmallVector<bool> packPaddings;
   for (int64_t packPadding : extractI64Array(getPackPaddings()))
     packPaddings.push_back(static_cast<bool>(packPadding));
 
   // Convert the padding values to attributes.
   SmallVector<Attribute> paddingValues;
   for (auto const &it :
        llvm::zip(getPaddingValues(), target->getOperandTypes())) {
-    Attribute attr = std::get<0>(it);
+    auto attr = std::get<0>(it).dyn_cast<TypedAttr>();
+    if (!attr) {
+      emitOpError("expects padding values to be typed attributes");
+      return DiagnosedSilenceableFailure::definiteFailure();
+    }
     Type elementType = getElementTypeOrSelf(std::get<1>(it));
     // Try to parse string attributes to obtain an attribute of element type.
     if (auto stringAttr = attr.dyn_cast<StringAttr>()) {
       paddingValues.push_back(
           parseAttribute(attr.cast<StringAttr>(), elementType));
       if (!paddingValues.back()) {
         auto diag = this->emitOpError("expects a padding that parses to ")
                     << elementType << ", got " << std::get<0>(it);

mlir/lib/Dialect/Linalg/Transforms/ElementwiseOpFusion.cpp

FoldScalarOrSplatConstant(MLIRContext *context, PatternBenefit benefit = 1)
       : OpRewritePattern<GenericOp>(context, benefit) {}
 
   LogicalResult matchAndRewrite(GenericOp genericOp,
                                 PatternRewriter &rewriter) const override {
     if (!genericOp.hasTensorSemantics())
       return failure();
     for (OpOperand *opOperand : genericOp.getInputOperands()) {
       Operation *def = opOperand->get().getDefiningOp();
-      Attribute constantAttr;
+      TypedAttr constantAttr;
       auto isScalarOrSplatConstantOp = [&constantAttr](Operation *def) -> bool {
         {
           DenseElementsAttr splatAttr;
           if (matchPattern(def, m_Constant<DenseElementsAttr>(&splatAttr)) &&
               splatAttr.isSplat() &&
               splatAttr.getType().getElementType().isIntOrFloat()) {
-            constantAttr = splatAttr.getSplatValue<Attribute>();
+            constantAttr = splatAttr.getSplatValue<TypedAttr>();
             return true;
           }
         }
         {
           IntegerAttr intAttr;
           if (matchPattern(def, m_Constant<IntegerAttr>(&intAttr))) {
             constantAttr = intAttr;
             return true;

mlir/lib/Dialect/Linalg/Transforms/Transforms.cpp

static FailureOr<Value> padOperandToSmallestStaticBoundingBox(
   });
   if (!nofold && hasStaticShape)
     return opOperand->get();
 
   // Fail if `paddingValues` specifies no padding value.
   if (opOperand->getOperandNumber() >= paddingValues.size())
     return failure();
   Attribute paddingAttr = paddingValues[opOperand->getOperandNumber()];
-  Value paddingValue = b.create<arith::ConstantOp>(
-      opToPad.getLoc(), paddingAttr.getType(), paddingAttr);
+  Type paddingType = b.getType<NoneType>();
+  if (auto typedAttr = paddingAttr.dyn_cast<TypedAttr>())
+    paddingType = typedAttr.getType();
+  Value paddingValue =
+      b.create<arith::ConstantOp>(opToPad.getLoc(), paddingType, paddingAttr);
 
   // Follow the use-def chain if `currOpOperand` is defined by a LinalgOp.
   OpOperand *currOpOperand = opOperand;
   while (auto linalgOp = currOpOperand->get().getDefiningOp<LinalgOp>()) {
     OpResult result = currOpOperand->get().cast<OpResult>();
     currOpOperand = linalgOp.getOutputOperand(result.getResultNumber());
   }
 

mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp

if (getInitialValue().has_value()) {
     Attribute initValue = getInitialValue().value();
     if (!initValue.isa<UnitAttr>() && !initValue.isa<ElementsAttr>())
       return emitOpError("initial value should be a unit or elements "
                          "attribute, but got ")
              << initValue;
 
     // Check that the type of the initial value is compatible with the type of
     // the global variable.
-    if (initValue.isa<ElementsAttr>()) {
-      Type initType = initValue.getType();
+    if (auto elementsAttr = initValue.dyn_cast<ElementsAttr>()) {
+      Type initType = elementsAttr.getType();
       Type tensorType = getTensorTypeFromMemRefType(memrefType);
       if (initType != tensorType)
         return emitOpError("initial value expected to be of type ")
                << tensorType << ", but was of type " << initType;
     }
   }
 
   if (Optional<uint64_t> alignAttr = getAlignment()) {

mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp

 using namespace mlir;
 
 //===----------------------------------------------------------------------===//
 // Common utility functions
 //===----------------------------------------------------------------------===//
 
 /// Returns the boolean value under the hood if the given `boolAttr` is a scalar
 /// or splat vector bool constant.
-static Optional<bool> getScalarOrSplatBoolAttr(Attribute boolAttr) {
-  if (!boolAttr)
+static Optional<bool> getScalarOrSplatBoolAttr(Attribute attr) {
+  if (!attr)
     return llvm::None;
 
-  auto type = boolAttr.getType();
-  if (type.isInteger(1)) {
-    auto attr = boolAttr.cast<BoolAttr>();
-    return attr.getValue();
-  }
-  if (auto vecType = type.cast<VectorType>()) {
-    if (vecType.getElementType().isInteger(1))
-      if (auto attr = boolAttr.dyn_cast<SplatElementsAttr>())
-        return attr.getSplatValue<bool>();
-  }
+  if (auto boolAttr = attr.dyn_cast<BoolAttr>())
+    return boolAttr.getValue();
+  if (auto splatAttr = attr.dyn_cast<SplatElementsAttr>())
+    if (splatAttr.getElementType().isInteger(1))
+      return splatAttr.getSplatValue<bool>();
   return llvm::None;
 }
 
 // Extracts an element from the given `composite` by following the given
 // `indices`. Returns a null Attribute if error happens.
 static Attribute extractCompositeElement(Attribute composite,
                                          ArrayRef<unsigned> indices) {
   // Check that given composite is a constant.

mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp

 //===----------------------------------------------------------------------===//
 
 ParseResult spirv::ConstantOp::parse(OpAsmParser &parser,
                                      OperationState &state) {
   Attribute value;
   if (parser.parseAttribute(value, kValueAttrName, state.attributes))
     return failure();
 
-  Type type = value.getType();
+  Type type = NoneType::get(parser.getContext());
+  if (auto typedAttr = value.dyn_cast<TypedAttr>())
+    type = typedAttr.getType();
   if (type.isa<NoneType, TensorType>()) {
     if (parser.parseColonType(type))
       return failure();
   }
 
   return parser.addTypeToList(type, state.types);
 }
 
 void spirv::ConstantOp::print(OpAsmPrinter &printer) {
   printer << ' ' << value();
   if (getType().isa<spirv::ArrayType>())
     printer << " : " << getType();
 }
 
 static LogicalResult verifyConstantType(spirv::ConstantOp op, Attribute value,
                                         Type opType) {
-  auto valueType = value.getType();
-
   if (value.isa<IntegerAttr, FloatAttr>()) {
+    auto valueType = value.cast<TypedAttr>().getType();
     if (valueType != opType)
       return op.emitOpError("result type (")
              << opType << ") does not match value type (" << valueType << ")";
     return success();
   }
   if (value.isa<DenseIntOrFPElementsAttr, SparseElementsAttr>()) {
+    auto valueType = value.cast<TypedAttr>().getType();
     if (valueType == opType)
       return success();
     auto arrayType = opType.dyn_cast<spirv::ArrayType>();
     auto shapedType = valueType.dyn_cast<ShapedType>();
     if (!arrayType)
       return op.emitOpError("result or element type (")
              << opType << ") does not match value type (" << valueType
              << "), must be the same or spv.array";
if (auto arrayAttr = value.dyn_cast<ArrayAttr>()) {
     Type elemType = arrayType.getElementType();
     for (Attribute element : arrayAttr.getValue()) {
       // Verify array elements recursively.
       if (failed(verifyConstantType(op, element, elemType)))
         return failure();
     }
     return success();
   }
-  return op.emitOpError("cannot have value of type ") << valueType;
+  return op.emitOpError("cannot have attribute: ") << value;
 }
 
 LogicalResult spirv::ConstantOp::verify() {
   // ODS already generates checks to make sure the result type is valid. We just
   // need to additionally check that the value's attribute type is consistent
   // with the result type.
   return verifyConstantType(*this, valueAttr(), getType());
 }

mlir/lib/Dialect/Vector/IR/VectorOps.cpp

 }
 
 OpFoldResult BroadcastOp::fold(ArrayRef<Attribute> operands) {
   if (getSourceType() == getVectorType())
     return getSource();
   if (!operands[0])
     return {};
   auto vectorType = getVectorType();
-  if (operands[0].getType().isIntOrIndexOrFloat())
+  if (operands[0].isa<IntegerAttr, FloatAttr>())
     return DenseElementsAttr::get(vectorType, operands[0]);
   if (auto attr = operands[0].dyn_cast<SplatElementsAttr>())
     return DenseElementsAttr::get(vectorType, attr.getSplatValue<Attribute>());
   return {};
 }
 
 namespace {
 
if (!getV1VectorType().isScalable() && !getV2VectorType().isScalable() &&
       isStepIndexArray(getMask(), getV1VectorType().getDimSize(0),
                        getV2VectorType().getDimSize(0)))
     return getV2();
 
   Attribute lhs = operands.front(), rhs = operands.back();
   if (!lhs || !rhs)
     return {};
 
-  auto lhsType = lhs.getType().cast<VectorType>();
+  auto lhsType = lhs.cast<DenseElementsAttr>().getType().cast<VectorType>();
   // Only support 1-D for now to avoid complicated n-D DenseElementsAttr
   // manipulation.
   if (lhsType.getRank() != 1)
     return {};
   int64_t lhsSize = lhsType.getDimSize(0);
 
   SmallVector<Attribute> results;
   auto lhsElements = lhs.cast<DenseElementsAttr>().getValues<Attribute>();

mlir/lib/IR/AsmPrinter.cpp

if (!attr) {
     os << "<<NULL ATTRIBUTE>>";
     return;
   }
 
   // Try to print an alias for this attribute.
   if (succeeded(printAlias(attr)))
     return;
 
-  auto attrType = attr.getType();
   if (!isa<BuiltinDialect>(attr.getDialect())) {
     printDialectAttribute(attr);
   } else if (auto opaqueAttr = attr.dyn_cast<OpaqueAttr>()) {
     printDialectSymbol(os, "#", opaqueAttr.getDialectNamespace(),
                        opaqueAttr.getAttrData());
   } else if (attr.isa<UnitAttr>()) {
     os << "unit";
     return;
   } else if (auto dictAttr = attr.dyn_cast<DictionaryAttr>()) {
     os << '{';
     interleaveComma(dictAttr.getValue(),
                     [&](NamedAttribute attr) { printNamedAttribute(attr); });
     os << '}';
 
   } else if (auto intAttr = attr.dyn_cast<IntegerAttr>()) {
-    if (attrType.isSignlessInteger(1)) {
+    Type intType = intAttr.getType();
+    if (intType.isSignlessInteger(1)) {
       os << (intAttr.getValue().getBoolValue() ? "true" : "false");
 
       // Boolean integer attributes always elides the type.
       return;
     }
 
     // Only print attributes as unsigned if they are explicitly unsigned or are
     // signless 1-bit values.  Indexes, signed values, and multi-bit signless
     // values print as signed.
     bool isUnsigned =
-        attrType.isUnsignedInteger() || attrType.isSignlessInteger(1);
+        intType.isUnsignedInteger() || intType.isSignlessInteger(1);
     intAttr.getValue().print(os, !isUnsigned);
 
     // IntegerAttr elides the type if I64.
-    if (typeElision == AttrTypeElision::May && attrType.isSignlessInteger(64))
+    if (typeElision == AttrTypeElision::May && intType.isSignlessInteger(64))
       return;
 
   } else if (auto floatAttr = attr.dyn_cast<FloatAttr>()) {
     printFloatValue(floatAttr.getValue(), os);
 
     // FloatAttr elides the type if F64.
-    if (typeElision == AttrTypeElision::May && attrType.isF64())
+    if (typeElision == AttrTypeElision::May && floatAttr.getType().isF64())
       return;
 
   } else if (auto strAttr = attr.dyn_cast<StringAttr>()) {
     printEscapedString(strAttr.getValue());
 
   } else if (auto arrayAttr = attr.dyn_cast<ArrayAttr>()) {
     os << '[';
     interleaveComma(arrayAttr.getValue(), [&](Attribute attr) {
case DenseArrayBaseAttr::EltType::I64:
       break;
     case DenseArrayBaseAttr::EltType::F32:
       os << "[:f32";
       break;
     case DenseArrayBaseAttr::EltType::F64:
       os << "[:f64";
       break;
     }
-    if (denseArrayAttr.getType().cast<ShapedType>().getRank())
+    if (denseArrayAttr.getType().getRank())
       os << " ";
     denseArrayAttr.printWithoutBraces(os);
     os << "]";
   } else if (auto locAttr = attr.dyn_cast<LocationAttr>()) {
     printLocation(locAttr);
   } else {
     llvm::report_fatal_error("Unknown builtin attribute");
   }
   // Don't print the type if we must elide it, or if it is a None type.
-  if (typeElision != AttrTypeElision::Must && !attrType.isa<NoneType>()) {
+  if (typeElision != AttrTypeElision::Must) {
+    if (auto typedAttr = attr.dyn_cast<TypedAttr>()) {
+      Type attrType = typedAttr.getType();
+      if (!attrType.isa<NoneType>()) {
         os << " : ";
         printType(attrType);
       }
     }
+  }
+}
 
 /// Print the integer element of a DenseElementsAttr.
 static void printDenseIntElement(const APInt &value, raw_ostream &os,
                                  bool isSigned) {
   if (value.getBitWidth() == 1)
     os << (value.getBoolValue() ? "true" : "false");
   else
     value.print(os, isSigned);

mlir/lib/IR/AttributeDetail.h

inline size_t getDenseElementBitWidth(Type eltType) {
   if (eltType.isIndex())
     return IndexType::kInternalStorageBitWidth;
   return eltType.getIntOrFloatBitWidth();
 }
 
 /// An attribute representing a reference to a dense vector or tensor object.
 struct DenseElementsAttributeStorage : public AttributeStorage {
 public:
-  DenseElementsAttributeStorage(ShapedType ty, bool isSplat)
-      : AttributeStorage(ty), isSplat(isSplat) {}
+  DenseElementsAttributeStorage(ShapedType type, bool isSplat)
+      : type(type), isSplat(isSplat) {}
 
+  ShapedType type;
   bool isSplat;
 };
 
 /// An attribute representing a reference to a dense vector or tensor object.
 struct DenseIntOrFPElementsAttrStorage : public DenseElementsAttributeStorage {
   DenseIntOrFPElementsAttrStorage(ShapedType ty, ArrayRef<char> data,
                                   bool isSplat = false)
       : DenseElementsAttributeStorage(ty, isSplat), data(data) {}
struct KeyTy {
     llvm::hash_code hashCode;
 
     /// A boolean that indicates if this data is a splat or not.
     bool isSplat;
   };
 
   /// Compare this storage instance with the provided key.
   bool operator==(const KeyTy &key) const {
-    if (key.type != getType())
+    if (key.type != type)
       return false;
 
     // For boolean splats we need to explicitly check that the first bit is the
     // same. Boolean values are packed at the bit level, and even though a splat
     // is detected the rest of the bits in the first byte may differ from the
     // splat value.
     if (key.type.getElementType().isInteger(1)) {
       if (key.isSplat != isSplat)
struct KeyTy {
     llvm::hash_code hashCode;
 
     /// A boolean that indicates if this data is a splat or not.
     bool isSplat;
   };
 
   /// Compare this storage instance with the provided key.
   bool operator==(const KeyTy &key) const {
-    if (key.type != getType())
+    if (key.type != type)
       return false;
 
     // Otherwise, we can default to just checking the data. StringRefs compare
     // by contents.
     return key.data == data;
   }
 
   /// Construct a key from a shaped type, StringRef data buffer, and a flag that

 };
 
 //===----------------------------------------------------------------------===//
 // StringAttr
 //===----------------------------------------------------------------------===//
 
 struct StringAttrStorage : public AttributeStorage {
   StringAttrStorage(StringRef value, Type type)
-      : AttributeStorage(type), value(value), referencedDialect(nullptr) {}
+      : type(type), value(value), referencedDialect(nullptr) {}
 
   /// The hash key is a tuple of the parameter types.
   using KeyTy = std::pair<StringRef, Type>;
   bool operator==(const KeyTy &key) const {
-    return value == key.first && getType() == key.second;
+    return value == key.first && type == key.second;
   }
   static ::llvm::hash_code hashKey(const KeyTy &key) {
     return DenseMapInfo<KeyTy>::getHashValue(key);
   }
 
   /// Define a construction method for creating a new instance of this
   /// storage.
   static StringAttrStorage *construct(AttributeStorageAllocator &allocator,
                                       const KeyTy &key) {
     return new (allocator.allocate<StringAttrStorage>())
         StringAttrStorage(allocator.copyInto(key.first), key.second);
   }
 
   /// Initialize the storage given an MLIRContext.
   void initialize(MLIRContext *context);
 
+  /// The type of the string.
+  Type type;
   /// The raw string value.
   StringRef value;
   /// If the string value contains a dialect namespace prefix (e.g.
   /// dialect.blah), this is the dialect referenced.
   Dialect *referencedDialect;
 };
 
 } // namespace detail
 } // namespace mlir
 
 #endif // ATTRIBUTEDETAIL_H_

mlir/lib/IR/BuiltinAttributeInterfaces.cpp

 //===----------------------------------------------------------------------===//
 
 #include "mlir/IR/BuiltinAttributeInterfaces.cpp.inc"
 
 //===----------------------------------------------------------------------===//
 // ElementsAttr
 //===----------------------------------------------------------------------===//
 
-ShapedType ElementsAttr::getType() const {
-  return Attribute::getType().cast<ShapedType>();
+Type ElementsAttr::getElementType(ElementsAttr elementsAttr) {
+  return elementsAttr.getType().getElementType();
 }
 
-Type ElementsAttr::getElementType(Attribute elementsAttr) {
-  return elementsAttr.getType().cast<ShapedType>().getElementType();
-}
-
-int64_t ElementsAttr::getNumElements(Attribute elementsAttr) {
-  return elementsAttr.getType().cast<ShapedType>().getNumElements();
+int64_t ElementsAttr::getNumElements(ElementsAttr elementsAttr) {
+  return elementsAttr.getType().getNumElements();
 }
 
 bool ElementsAttr::isValidIndex(ShapedType type, ArrayRef<uint64_t> index) {
   // Verify that the rank of the indices matches the held type.
   int64_t rank = type.getRank();
   if (rank == 0 && index.size() == 1 && index[0] == 0)
     return true;
   if (rank != static_cast<int64_t>(index.size()))
     return false;
 
   // Verify that all of the indices are within the shape dimensions.
   ArrayRef<int64_t> shape = type.getShape();
   return llvm::all_of(llvm::seq<int>(0, rank), [&](int i) {
     int64_t dim = static_cast<int64_t>(index[i]);
     return 0 <= dim && dim < shape[i];
   });
 }
-bool ElementsAttr::isValidIndex(Attribute elementsAttr,
+bool ElementsAttr::isValidIndex(ElementsAttr elementsAttr,
                                 ArrayRef<uint64_t> index) {
-  return isValidIndex(elementsAttr.getType().cast<ShapedType>(), index);
+  return isValidIndex(elementsAttr.getType(), index);
 }
 
 uint64_t ElementsAttr::getFlattenedIndex(Type type, ArrayRef<uint64_t> index) {
   ShapedType shapeType = type.cast<ShapedType>();
   assert(isValidIndex(shapeType, index) &&
          "expected valid multi-dimensional index");
 
   // Reduce the provided multidimensional index into a flattended 1D row-major

mlir/lib/IR/BuiltinAttributes.cpp

 /// Twine support for StringAttr.
 StringAttr StringAttr::get(const Twine &twine, Type type) {
   SmallVector<char, 32> tempStr;
   return Base::get(type.getContext(), twine.toStringRef(tempStr), type);
 }
 
 StringRef StringAttr::getValue() const { return getImpl()->value; }
 
+Type StringAttr::getType() const { return getImpl()->type; }
+
 Dialect *StringAttr::getReferencedDialect() const {
   return getImpl()->referencedDialect;
 }
 
 //===----------------------------------------------------------------------===//
 // FloatAttr
 //===----------------------------------------------------------------------===//
 

 /// Custom storage to ensure proper memory alignment for the allocation of
 /// DenseArray of any element type.
 struct mlir::detail::DenseArrayBaseAttrStorage : public AttributeStorage {
   using KeyTy = std::tuple<ShapedType, DenseArrayBaseAttr::EltType,
                            ::llvm::ArrayRef<char>>;
   DenseArrayBaseAttrStorage(ShapedType type,
                             DenseArrayBaseAttr::EltType eltType,
                             ::llvm::ArrayRef<char> elements)
-      : AttributeStorage(type), eltType(eltType), elements(elements) {}
+      : type(type), eltType(eltType), elements(elements) {}
 
   bool operator==(const KeyTy &tblgenKey) const {
-    return (getType() == std::get<0>(tblgenKey)) &&
+    return (type == std::get<0>(tblgenKey)) &&
            (eltType == std::get<1>(tblgenKey)) &&
            (elements == std::get<2>(tblgenKey));
   }
 
-  static ::llvm::hash_code hashKey(const KeyTy &tblgenKey) {
-    return ::llvm::hash_combine(std::get<0>(tblgenKey), std::get<1>(tblgenKey),
-                                std::get<2>(tblgenKey));
+  static llvm::hash_code hashKey(const KeyTy &tblgenKey) {
+    return llvm::hash_combine(std::get<0>(tblgenKey), std::get<1>(tblgenKey),
   }
 
   static DenseArrayBaseAttrStorage *
   construct(AttributeStorageAllocator &allocator, const KeyTy &tblgenKey) {
     auto type = std::get<0>(tblgenKey);
     auto eltType = std::get<1>(tblgenKey);
     auto elements = std::get<2>(tblgenKey);
     if (!elements.empty()) {
       char *alloc = static_cast<char *>(
           allocator.allocate(elements.size(), alignof(uint64_t)));
       std::uninitialized_copy(elements.begin(), elements.end(), alloc);
       elements = ArrayRef<char>(alloc, elements.size());
     }
     return new (allocator.allocate<DenseArrayBaseAttrStorage>())
         DenseArrayBaseAttrStorage(type, eltType, elements);
   }
 
+  ShapedType type;
   DenseArrayBaseAttr::EltType eltType;
-  ::llvm::ArrayRef<char> elements;
+  ArrayRef<char> elements;
 };
 
 DenseArrayBaseAttr::EltType DenseArrayBaseAttr::getElementType() const {
   return getImpl()->eltType;
 }
 
+ShapedType DenseArrayBaseAttr::getType() const { return getImpl()->type; }
+
 const int8_t *
 DenseArrayBaseAttr::value_begin_impl(OverloadToken<int8_t>) const {
   return cast<DenseI8ArrayAttr>().asArrayRef().begin();
 }
 const int16_t *
 DenseArrayBaseAttr::value_begin_impl(OverloadToken<int16_t>) const {
   return cast<DenseI16ArrayAttr>().asArrayRef().begin();
 }

 }
 
 DenseElementsAttr DenseElementsAttr::get(ShapedType type,
                                          ArrayRef<Attribute> values) {
   assert(hasSameElementsOrSplat(type, values));
 
   // If the element type is not based on int/float/index, assume it is a string
   // type.
-  auto eltType = type.getElementType();
-  if (!type.getElementType().isIntOrIndexOrFloat()) {
+  Type eltType = type.getElementType();
+  if (!eltType.isIntOrIndexOrFloat()) {
     SmallVector<StringRef, 8> stringValues;
     stringValues.reserve(values.size());
     for (Attribute attr : values) {
       assert(attr.isa<StringAttr>() &&
              "expected string value for non integer/index/float element");
       stringValues.push_back(attr.cast<StringAttr>().getValue());
     }
     return get(type, stringValues);
   }
 
   // Otherwise, get the raw storage width to use for the allocation.
   size_t bitWidth = getDenseElementBitWidth(eltType);
   size_t storageBitWidth = getDenseElementStorageWidth(bitWidth);
 
   // Compress the attribute values into a character buffer.
   SmallVector<char, 8> data(
       llvm::divideCeil(storageBitWidth * values.size(), CHAR_BIT));
   APInt intVal;
   for (unsigned i = 0, e = values.size(); i < e; ++i) {
-    assert(eltType == values[i].getType() &&
-           "expected attribute value to have element type");
-    if (eltType.isa<FloatType>())
-      intVal = values[i].cast<FloatAttr>().getValue().bitcastToAPInt();
-    else if (eltType.isa<IntegerType, IndexType>())
-      intVal = values[i].cast<IntegerAttr>().getValue();
-    else
-      llvm_unreachable("unexpected element type");
+    if (auto floatAttr = values[i].dyn_cast<FloatAttr>()) {
+      assert(floatAttr.getType() == eltType &&
+             "expected float attribute type to equal element type");
+      intVal = floatAttr.getValue().bitcastToAPInt();
+    } else {
+      auto intAttr = values[i].cast<IntegerAttr>();
+      assert(intAttr.getType() == eltType &&
+             "expected integer attribute type to equal element type");
+      intVal = intAttr.getValue();
+    }
 
     assert(intVal.getBitWidth() == bitWidth &&
            "expected value to have same bitwidth as element type");
     writeBits(data.data(), i * storageBitWidth, intVal);
   }
 
   // Handle the special encoding of splat of bool.
-  if (values.size() == 1 && values[0].getType().isInteger(1))
+  if (values.size() == 1 && eltType.isInteger(1))
     data[0] = data[0] ? -1 : 0;
 
   return DenseIntOrFPElementsAttr::getRaw(type, data);
 }
 
 DenseElementsAttr DenseElementsAttr::get(ShapedType type,
                                          ArrayRef<bool> values) {
   assert(hasSameElementsOrSplat(type, values));

 }
 
 DenseElementsAttr DenseElementsAttr::mapValues(
     Type newElementType, function_ref<APInt(const APFloat &)> mapping) const {
   return cast<DenseFPElementsAttr>().mapValues(newElementType, mapping);
 }
 
 ShapedType DenseElementsAttr::getType() const {
-  return Attribute::getType().cast<ShapedType>();
+  return static_cast<const DenseElementsAttributeStorage *>(impl)->type;
 }
 
 Type DenseElementsAttr::getElementType() const {
   return getType().getElementType();
 }
 
 int64_t DenseElementsAttr::getNumElements() const {
   return getType().getNumElements();
DenseElementsAttr DenseFPElementsAttr::mapValues(
   auto newArrayType =
       mappingHelper(mapping, *this, getType(), newElementType, elementData);
 
   return getRaw(newArrayType, elementData);
 }
 
 /// Method for supporting type inquiry through isa, cast and dyn_cast.
 bool DenseFPElementsAttr::classof(Attribute attr) {
-  return attr.isa<DenseElementsAttr>() &&
-         attr.getType().cast<ShapedType>().getElementType().isa<FloatType>();
+  if (auto denseAttr = attr.dyn_cast<DenseElementsAttr>())
+    return denseAttr.getType().getElementType().isa<FloatType>();
+  return false;
 }
 
 //===----------------------------------------------------------------------===//
 // DenseIntElementsAttr
 //===----------------------------------------------------------------------===//
 
 DenseElementsAttr DenseIntElementsAttr::mapValues(
     Type newElementType, function_ref<APInt(const APInt &)> mapping) const {
   llvm::SmallVector<char, 8> elementData;
   auto newArrayType =
       mappingHelper(mapping, *this, getType(), newElementType, elementData);
   return getRaw(newArrayType, elementData);
 }
 
 /// Method for supporting type inquiry through isa, cast and dyn_cast.
 bool DenseIntElementsAttr::classof(Attribute attr) {
-  return attr.isa<DenseElementsAttr>() &&
-         attr.getType().cast<ShapedType>().getElementType().isIntOrIndex();
+  if (auto denseAttr = attr.dyn_cast<DenseElementsAttr>())
+    return denseAttr.getType().getElementType().isIntOrIndex();
+  return false;
 }
 
 //===----------------------------------------------------------------------===//
 // OpaqueElementsAttr
 //===----------------------------------------------------------------------===//
 
 bool OpaqueElementsAttr::decode(ElementsAttr &result) {
   Dialect *dialect = getContext()->getLoadedDialect(getDialect());

mlir/lib/IR/MLIRContext.cpp

StorageUniquer &MLIRContext::getAttributeUniquer() {
   return getImpl().attributeUniquer;
 }
 
 /// Initialize the given attribute storage instance.
 void AttributeUniquer::initializeAttributeStorage(AttributeStorage *storage,
                                                   MLIRContext *ctx,
                                                   TypeID attrID) {
   storage->initializeAbstractAttribute(AbstractAttribute::lookup(attrID, ctx));
-
-  // If the attribute did not provide a type, then default to NoneType.
-  if (!storage->getType())
-    storage->setType(NoneType::get(ctx));
 }
 
 BoolAttr BoolAttr::get(MLIRContext *context, bool value) {
   return value ? context->getImpl().trueAttr : context->getImpl().falseAttr;
 }
 
 UnitAttr UnitAttr::get(MLIRContext *context) {
   return context->getImpl().unitAttr;

mlir/lib/IR/TypeUtilities.cpp

Type mlir::getElementTypeOrSelf(Type type) {
   return type;
 }
 
 Type mlir::getElementTypeOrSelf(Value val) {
   return getElementTypeOrSelf(val.getType());
 }
 
 Type mlir::getElementTypeOrSelf(Attribute attr) {
-  return getElementTypeOrSelf(attr.getType());
+  if (auto typedAttr = attr.dyn_cast<TypedAttr>())
+    return getElementTypeOrSelf(typedAttr.getType());
+  return {};
 }
 
 SmallVector<Type, 10> mlir::getFlattenedTypes(TupleType t) {
   SmallVector<Type, 10> fTypes;
   t.getFlattenedTypes(fTypes);
   return fTypes;
 }
 

mlir/lib/Parser/DialectSymbolParser.cpp

Attribute attr = parseExtendedSymbol<Attribute>(
 
         // Otherwise, form a new opaque attribute.
         return OpaqueAttr::getChecked(
             [&] { return emitError(loc); }, StringAttr::get(ctx, dialectName),
             symbolData, attrType ? attrType : NoneType::get(ctx));
       });
 
   // Ensure that the attribute has the same type as requested.
-  if (attr && type && attr.getType() != type) {
+  if (attr && type) {
+    auto typedAttr = attr.dyn_cast<TypedAttr>();
+    if (typedAttr && typedAttr.getType() != type) {
       emitError("attribute type different than expected: expected ")
-        << type << ", but got " << attr.getType();
+          << type << ", but got " << typedAttr.getType();
       return nullptr;
     }
+  }
   return attr;
 }
 
 /// Parse an extended type.
 ///
 ///   extended-type ::= (dialect-type | type-alias)
 ///   dialect-type  ::= `!` dialect-namespace `<` `"` type-data `"` `>`
 ///   dialect-type  ::= `!` alias-name pretty-dialect-attribute-body?

mlir/lib/Rewrite/ByteCode.cpp

LLVM_DEBUG(llvm::dbgs() << "  * Operation: " << *op << "\n"
                           << "  * Result: " << attr << "\n");
   memory[memIndex] = attr.getAsOpaquePointer();
 }
 
 void ByteCodeExecutor::executeGetAttributeType() {
   LLVM_DEBUG(llvm::dbgs() << "Executing GetAttributeType:\n");
   unsigned memIndex = read();
   Attribute attr = read<Attribute>();
-  Type type = attr ? attr.getType() : Type();
+  Type type;
+  if (auto typedAttr = attr.dyn_cast<TypedAttr>())
+    type = typedAttr.getType();
 
   LLVM_DEBUG(llvm::dbgs() << "  * Attribute: " << attr << "\n"
                           << "  * Result: " << type << "\n");
   memory[memIndex] = type.getAsOpaquePointer();
 }
 
 void ByteCodeExecutor::executeGetDefiningOp() {
   LLVM_DEBUG(llvm::dbgs() << "Executing GetDefiningOp:\n");

mlir/lib/TableGen/AttrOrTypeDef.cpp

 
 bool AttrOrTypeParameter::isOptional() const {
   // Parameters with default values are automatically optional.
   return getDefValue<llvm::BitInit>("isOptional").value_or(false) ||
          getDefaultValue();
 }
 
 Optional<StringRef> AttrOrTypeParameter::getDefaultValue() const {
-  return getDefValue<llvm::StringInit>("defaultValue");
+  Optional<StringRef> result = getDefValue<llvm::StringInit>("defaultValue");
+  return result && !result->empty() ? result : llvm::None;
 }
 
 llvm::Init *AttrOrTypeParameter::getDef() const { return def->getArg(index); }
 
 //===----------------------------------------------------------------------===//
 // AttributeSelfTypeParameter
 //===----------------------------------------------------------------------===//
 
 bool AttributeSelfTypeParameter::classof(const AttrOrTypeParameter *param) {
   llvm::Init *paramDef = param->getDef();
   if (auto *paramDefInit = dyn_cast<llvm::DefInit>(paramDef))
     return paramDefInit->getDef()->isSubClassOf("AttributeSelfTypeParameter");
   return false;
 }

mlir/lib/Target/Cpp/TranslateToCpp.cpp

if (auto sAttr = attr.dyn_cast<SymbolRefAttr>()) {
     os << sAttr.getRootReference().getValue();
     return success();
   }
 
   // Print type attributes.
   if (auto type = attr.dyn_cast<TypeAttr>())
     return emitType(loc, type.getValue());
 
-  return emitError(loc, "cannot emit attribute of type ") << attr.getType();
+  return emitError(loc, "cannot emit attribute: ") << attr;
 }
 
 LogicalResult CppEmitter::emitOperands(Operation &op) {
   auto emitOperandName = [&](Value result) -> LogicalResult {
     if (!hasValueInScope(result))
       return op.emitOpError() << "operand value not in scope";
     os << getOrCreateName(result);
     return success();

mlir/lib/Target/SPIRV/Serialization/Serializer.cpp

if (!isSpec) {
     // We can de-duplicate normal constants, but not specialization constants.
     if (auto id = getConstantID(boolAttr)) {
       return id;
     }
   }
 
   // Process the type for this bool literal
   uint32_t typeID = 0;
-  if (failed(processType(loc, boolAttr.getType(), typeID))) {
+  if (failed(
+          processType(loc, boolAttr.cast<IntegerAttr>().getType(), typeID))) {
     return 0;
   }
 
   auto resultID = getNextID();
   auto opcode = boolAttr.getValue()
                     ? (isSpec ? spirv::Opcode::OpSpecConstantTrue
                               : spirv::Opcode::OpConstantTrue)
                     : (isSpec ? spirv::Opcode::OpSpecConstantFalse

mlir/test/Dialect/LLVMIR/invalid.mlir

llvm.func @array_attribute_one_element() -> !llvm.struct<(f64, f64)> {
   // expected-error @+1 {{expected array attribute with two elements, representing a complex constant}}
   %0 = llvm.mlir.constant([1.0 : f64]) : !llvm.struct<(f64, f64)>
   llvm.return %0 : !llvm.struct<(f64, f64)>
 }
 
 // -----
 
 llvm.func @array_attribute_two_different_types() -> !llvm.struct<(f64, f64)> {
-  // expected-error @+1 {{expected array attribute with two elements, representing a complex constant}}
+  // expected-error @+1 {{expected array attribute with two elements of the same type}}
   %0 = llvm.mlir.constant([1.0 : f64, 1.0 : f32]) : !llvm.struct<(f64, f64)>
   llvm.return %0 : !llvm.struct<(f64, f64)>
 }
 
 // -----
 
 llvm.func @struct_wrong_attribute_type() -> !llvm.struct<(f64, f64)> {
   // expected-error @+1 {{expected array attribute with two elements, representing a complex constant}}

 
 // -----
 
 func.func @nvvm_invalid_mma_0(%a0 : f16, %a1 : f16,
                          %b0 : vector<2xf16>, %b1 : vector<2xf16>,
                          %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32,
                          %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) {
   // expected-error@+1 {{Could not match types for the A operands; expected one of 2xvector<2xf16> but got f16, f16}}
-  %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7] 
+  %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7]
     {layoutA=#nvvm.mma_layout<row>, layoutB=#nvvm.mma_layout<col>, shape = #nvvm.shape<m = 8, n = 8, k = 4>} : (f16, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>
   llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>
 }
 
 // -----
 
 func.func @nvvm_invalid_mma_1(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
                          %b0 : vector<2xf16>, %b1 : vector<2xf16>,
                          %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32,
                          %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) {
   // expected-error@+1 {{Could not match allowed types for the result; expected one of !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>, !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> but got !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f16)>}}
   %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7]
     {layoutA=#nvvm.mma_layout<row>, layoutB=#nvvm.mma_layout<col>, shape = #nvvm.shape<m = 8, n = 8, k = 4>} : (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f16)>
   llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f16)>
 }
 
 // -----
 
 func.func @nvvm_invalid_mma_2(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
                          %b0 : vector<2xf16>, %b1 : vector<2xf16>,
                          %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32,
                          %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) {
   // expected-error@+1 {{op requires attribute 'layoutA'}}
-  %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7] 
+  %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7]
     {shape = #nvvm.shape<m = 8, n = 8, k = 4>}: (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>
   llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>
 }
 
 // -----
 
 func.func @nvvm_invalid_mma_3(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
                          %b0 : vector<2xf16>, %b1 : vector<2xf16>,
                          %c0 : vector<2xf16>, %c1 : vector<2xf16>) {
   // expected-error@+1 {{unimplemented variant for MMA shape <8, 8, 16>}}
   %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1] {layoutA=#nvvm.mma_layout<row>, layoutB=#nvvm.mma_layout<col>, shape = #nvvm.shape<m = 8, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
   llvm.return %0 : !llvm.struct<(vector<2xf16>, vector<2xf16>)>
 }
 
 // -----
 
 func.func @nvvm_invalid_mma_8(%a0 : i32, %a1 : i32,
                                %b0 : i32,
                                %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) {
   // expected-error@+1 {{op requires b1Op attribute}}
   %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]
     {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,
-     multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>,     
+     multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>,
      shape = #nvvm.shape<m = 16, n = 8, k = 128>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>
   llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>
 }
 
 // -----
 
 func.func @atomicrmw_expected_ptr(%f32 : f32) {
   // expected-error@+1 {{operand #0 must be LLVM pointer to floating point LLVM type or integer}}

mlir/test/Dialect/SPIRV/IR/structure-ops.mlir

func.func @const() -> () {
   %8 = spv.Constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
   %9 = spv.Constant [[dense<3.0> : vector<2xf32>]] : !spv.array<1 x !spv.array<1xvector<2xf32>>>
   return
 }
 
 // -----
 
 func.func @unaccepted_std_attr() -> () {
-  // expected-error @+1 {{cannot have value of type 'none'}}
+  // expected-error @+1 {{cannot have attribute: unit}}
   %0 = spv.Constant unit : none
   return
 }
 
 // -----
 
 func.func @array_constant() -> () {
   // expected-error @+1 {{result or element type ('vector<2xf32>') does not match value type ('vector<2xi32>')}}

mlir/test/IR/file-metadata-resources.mlir

 // RUN: mlir-opt %s -split-input-file | FileCheck %s
 
 // Check that we only preserve the blob that got referenced.
 // CHECK:      test: {
 // CHECK-NEXT:   blob1: "0x08000000010000000000000002000000000000000300000000000000"
 // CHECK-NEXT: }
 
-module attributes { test.blob_ref = #test.e1di64_elements<blob1> } {}
+module attributes { test.blob_ref = #test.e1di64_elements<blob1> : tensor<*xi1>} {}
 
 {-#
   dialect_resources: {
     test: {
       blob1: "0x08000000010000000000000002000000000000000300000000000000",
       blob2: "0x08000000040000000000000005000000000000000600000000000000"
     }
   }
 #-}

mlir/test/lib/Dialect/Test/TestAttrDefs.td

 }
 def CompoundAttrNested : Test_Attr<"CompoundAttrNested"> {
   let mnemonic = "cmpnd_nested";
   let parameters = (ins CompoundAttrA : $nested );
   let assemblyFormat = "`<` `nested` `=` $nested `>`";
 }
 
 // An attribute testing AttributeSelfTypeParameter.
-def AttrWithSelfTypeParam : Test_Attr<"AttrWithSelfTypeParam"> {
+def AttrWithSelfTypeParam
+    : Test_Attr<"AttrWithSelfTypeParam", [TypedAttrInterface]> {
   let mnemonic = "attr_with_self_type_param";
   let parameters = (ins AttributeSelfTypeParameter<"">:$type);
   let assemblyFormat = "";
 }
 
 // An attribute testing AttributeSelfTypeParameter.
-def AttrWithTypeBuilder : Test_Attr<"AttrWithTypeBuilder"> {
+def AttrWithTypeBuilder
+    : Test_Attr<"AttrWithTypeBuilder", [TypedAttrInterface]> {
   let mnemonic = "attr_with_type_builder";
-  let parameters = (ins "::mlir::IntegerAttr":$attr);
-  let typeBuilder = "$_attr.getType()";
-  let hasCustomAssemblyFormat = 1;
+  let parameters = (ins
+    "::mlir::IntegerAttr":$attr,
+    AttributeSelfTypeParameter<"", "mlir::Type", "$attr.getType()">:$type
+  );
+  let assemblyFormat = "$attr";
 }
 
 def TestAttrTrait : NativeAttrTrait<"TestAttrTrait">;
 
 // The definition of a singleton attribute that has a trait.
 def AttrWithTrait : Test_Attr<"AttrWithTrait", [TestAttrTrait]> {
   let mnemonic = "attr_with_trait";
 }
 
 // Test support for ElementsAttrInterface.
 def TestI64ElementsAttr : Test_Attr<"TestI64Elements", [
-    ElementsAttrInterface
+    ElementsAttrInterface, TypedAttrInterface
   ]> {
   let mnemonic = "i64_elements";
   let parameters = (ins
     AttributeSelfTypeParameter<"", "::mlir::ShapedType">:$type,
     ArrayRefParameter<"uint64_t">:$elements
   );
   let extraClassDeclaration = [{
     /// The set of data types that can be iterated by this attribute.
def TestAttrWithTypeParam : Test_Attr<"TestAttrWithTypeParam"> {
   let parameters = (ins "::mlir::IntegerType":$int_type,
                         "::mlir::Type":$any_type);
   let mnemonic = "attr_with_type";
   let assemblyFormat = "`<` $int_type `,` $any_type `>`";
 }
 
 // Test self type parameter with assembly format.
 def TestAttrSelfTypeParameterFormat
-    : Test_Attr<"TestAttrSelfTypeParameterFormat"> {
+    : Test_Attr<"TestAttrSelfTypeParameterFormat", [TypedAttrInterface]> {
   let parameters = (ins "int":$a, AttributeSelfTypeParameter<"">:$type);
 
   let mnemonic = "attr_self_type_format";
   let assemblyFormat = "`<` $a `>`";
 }
 
 // Test overridding attribute builders with a custom builder.
 def TestOverrideBuilderAttr : Test_Attr<"TestOverrideBuilder"> {
   let mnemonic = "override_builder";
   let parameters = (ins "int":$a);
   let assemblyFormat = "`<` $a `>`";
 
   let skipDefaultBuilders = 1;
   let genVerifyDecl = 1;
   let builders = [AttrBuilder<(ins "int":$a), [{
     return ::mlir::IntegerAttr::get(::mlir::IndexType::get($_ctxt), a);
   }], "::mlir::Attribute">];
 }
 
 // Test simple extern 1D vector using ElementsAttrInterface.
 def TestExtern1DI64ElementsAttr : Test_Attr<"TestExtern1DI64Elements", [
-    ElementsAttrInterface
+    ElementsAttrInterface, TypedAttrInterface
   ]> {
   let mnemonic = "e1di64_elements";
   let parameters = (ins
     AttributeSelfTypeParameter<"", "::mlir::ShapedType">:$type,
     ResourceHandleParameter<"TestExternalElementsDataHandle">:$handle
   );
   let extraClassDeclaration = [{
     /// Return the elements referenced by this attribute.

mlir/test/lib/Dialect/Test/TestAttributes.cpp

 #include "llvm/ADT/TypeSwitch.h"
 #include "llvm/ADT/bit.h"
 #include "llvm/Support/ErrorHandling.h"
 
 using namespace mlir;
 using namespace test;
 
 //===----------------------------------------------------------------------===//
-// AttrWithTypeBuilderAttr
-//===----------------------------------------------------------------------===//
-
-Attribute AttrWithTypeBuilderAttr::parse(AsmParser &parser, Type type) {
-  IntegerAttr element;
-  if (parser.parseAttribute(element))
-    return Attribute();
-  return get(parser.getContext(), element);
-}
-
-void AttrWithTypeBuilderAttr::print(AsmPrinter &printer) const {
-  printer << " " << getAttr();
-}
-
-//===----------------------------------------------------------------------===//
 // CompoundAAttr
 //===----------------------------------------------------------------------===//
 
 Attribute CompoundAAttr::parse(AsmParser &parser, Type type) {
   int widthOfSomething;
   Type oneType;
   SmallVector<int, 4> arrayOfInts;
   if (parser.parseLess() || parser.parseInteger(widthOfSomething) ||
return emitError()
            << elements.size() << ", but expected: " << type.getNumElements();
   }
   if (type.getRank() != 1 || !type.getElementType().isSignlessInteger(64))
     return emitError() << "expected single rank 64-bit shape type, but got: "
                        << type;
   return success();
 }
 
-LogicalResult TestAttrWithFormatAttr::verify(
-    function_ref<InFlightDiagnostic()> emitError, int64_t one, std::string two,
-    IntegerAttr three, ArrayRef<int> four,
-    ArrayRef<AttrWithTypeBuilderAttr> arrayOfAttrWithTypeBuilderAttr) {
+LogicalResult
+TestAttrWithFormatAttr::verify(function_ref<InFlightDiagnostic()> emitError,
+                               int64_t one, std::string two, IntegerAttr three,
+                               ArrayRef<int> four,
+                               ArrayRef<AttrWithTypeBuilderAttr> arrayOfAttrs) {
   if (four.size() != static_cast<unsigned>(one))
     return emitError() << "expected 'one' to equal 'four.size()'";
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Utility Functions for Generated Attributes
 //===----------------------------------------------------------------------===//

mlir/test/lib/Dialect/Test/TestOps.td

 def OperandsHaveSameType :
     TEST_Op<"operands_have_same_type", [AllTypesMatch<["x", "y"]>]> {
   let arguments = (ins AnyType:$x, AnyType:$y);
 }
 
 def ResultHasSameTypeAsAttr :
     TEST_Op<"result_has_same_type_as_attr",
             [AllTypesMatch<["attr", "result"]>]> {
-  let arguments = (ins AnyAttr:$attr);
+  let arguments = (ins TypedAttrInterface:$attr);
   let results = (outs AnyType:$result);
   let assemblyFormat = "$attr `->` type($result) attr-dict";
 }
 
 def OperandZeroAndResultHaveSameType :
     TEST_Op<"operand0_and_result_have_same_type",
             [AllTypesMatch<["x", "res"]>]> {
   let arguments = (ins AnyType:$x, AnyType:$y);
def FormatAllTypesMatchVarOp : TEST_Op<"format_all_types_match_var", [
   let arguments = (ins AnyType:$value1, AnyType:$value2);
   let results = (outs AnyType:$result);
   let assemblyFormat = "attr-dict $value1 `,` $value2 `:` type($value1)";
 }
 
 def FormatAllTypesMatchAttrOp : TEST_Op<"format_all_types_match_attr", [
     AllTypesMatch<["value1", "value2", "result"]>
   ]> {
-  let arguments = (ins AnyAttr:$value1, AnyType:$value2);
+  let arguments = (ins TypedAttrInterface:$value1, AnyType:$value2);
   let results = (outs AnyType:$result);
   let assemblyFormat = "attr-dict $value1 `,` $value2";
 }
 
 //===----------------------------------------------------------------------===//
 // TypesMatchWith type inference
 
 def FormatTypesMatchVarOp : TEST_Op<"format_types_match_var", [
def FormatTypesMatchVariadicOp : TEST_Op<"format_types_match_variadic", [
   let arguments = (ins Variadic<AnyType>:$value);
   let results = (outs Variadic<AnyType>:$result);
   let assemblyFormat = "attr-dict $value `:` type($value)";
 }
 
 def FormatTypesMatchAttrOp : TEST_Op<"format_types_match_attr", [
     TypesMatchWith<"result type matches constant", "value", "result", "$_self">
   ]> {
-  let arguments = (ins AnyAttr:$value);
+  let arguments = (ins TypedAttrInterface:$value);
   let results = (outs AnyType:$result);
   let assemblyFormat = "attr-dict $value";
 }
 
 def FormatTypesMatchContextOp : TEST_Op<"format_types_match_context", [
     TypesMatchWith<"tuple result type matches operand type", "value", "result",
         "::mlir::TupleType::get($_ctxt, $_self)">
   ]> {

mlir/test/mlir-tblgen/attr-or-type-format.td

def AttrC : TestAttr<"TestF"> {
   let parameters = (ins "int":$v0, "int":$v1);
 
   let mnemonic = "attr_c";
   let assemblyFormat = "params";
 }
 
 /// Test attribute with self type parameter
 
-// ATTR: TestGAttr::parse
-// ATTR:   return TestGAttr::get
-// ATTR:     odsType
+// ATTR-LABEL: Attribute TestGAttr::parse
+// ATTR: if (odsType)
+// ATTR:   if (auto reqType = odsType.dyn_cast<::mlir::Type>())
+// ATTR:     _result_type = reqType
+// ATTR: TestGAttr::get
+// ATTR-NEXT: *_result_a
+// ATTR-NEXT: _result_type.value_or(::mlir::NoneType::get(
 def AttrD : TestAttr<"TestG"> {
   let parameters = (ins "int":$a, AttributeSelfTypeParameter<"">:$type);
   let mnemonic = "attr_d";
   let assemblyFormat = "$a";
 }
 
 /// Test type parser and printer that mix variables and struct are generated
 /// correctly.

mlir/test/mlir-tblgen/attrdefs.td

 // DECL:   return {"cmpnd_a"};
 // DECL: }
 // DECL: static ::mlir::Attribute parse(
 // DECL-SAME: ::mlir::AsmParser &odsParser, ::mlir::Type odsType);
 // DECL: void print(::mlir::AsmPrinter &odsPrinter) const;
 // DECL: int getWidthOfSomething() const;
 // DECL: ::test::SimpleTypeA getExampleTdType() const;
 // DECL: ::llvm::APFloat getApFloat() const;
+// DECL: ::mlir::Type getInner() const;
 
 // Check that AttributeSelfTypeParameter is handled properly.
 // DEF-LABEL: struct CompoundAAttrStorage
 // DEF: CompoundAAttrStorage(
-// DEF-SAME: : ::mlir::AttributeStorage(inner),
+// DEF-SAME: inner(inner)
 
 // DEF: bool operator==(const KeyTy &tblgenKey) const {
 // DEF-NEXT: return
 // DEF-SAME: (widthOfSomething == std::get<0>(tblgenKey)) &&
 // DEF-SAME: (exampleTdType == std::get<1>(tblgenKey)) &&
 // DEF-SAME: (apFloat.bitwiseIsEqual(std::get<2>(tblgenKey))) &&
 // DEF-SAME: (dims == std::get<3>(tblgenKey)) &&
-// DEF-SAME: (getType() == std::get<4>(tblgenKey));
+// DEF-SAME: (inner == std::get<4>(tblgenKey));
 
 // DEF: static CompoundAAttrStorage *construct
 // DEF: return new (allocator.allocate<CompoundAAttrStorage>())
 // DEF-SAME: CompoundAAttrStorage(widthOfSomething, exampleTdType, apFloat, dims, inner);
 
 // DEF: ::mlir::Type CompoundAAttr::getInner() const {
-// DEF-NEXT: return getImpl()->getType().cast<::mlir::Type>();
+// DEF-NEXT: return getImpl()->inner;
 }
 
 def C_IndexAttr : TestAttr<"Index"> {
     let mnemonic = "index";
 
     let parameters = (
       ins
       StringRefParameter<"Label for index">:$label
let parameters = (
     ins
     "int": $num
   );
 // DECL-LABEL: struct SingleParameterAttrStorage;
 // DECL-LABEL: class SingleParameterAttr
 // DECL-SAME:  detail::SingleParameterAttrStorage
 }
 
-// An attribute testing AttributeSelfTypeParameter.
-def E_AttrWithTypeBuilder : TestAttr<"AttrWithTypeBuilder"> {
-  let mnemonic = "attr_with_type_builder";
-  let parameters = (ins "::mlir::IntegerAttr":$attr);
-  let typeBuilder = "$_attr.getType()";
-  let hasCustomAssemblyFormat = 1;
-}
-
-// DEF-LABEL: struct AttrWithTypeBuilderAttrStorage
-// DEF: AttrWithTypeBuilderAttrStorage(::mlir::IntegerAttr attr)
-// DEF-SAME: : ::mlir::AttributeStorage(attr.getType()), attr(attr)
-
 def F_ParamWithAccessorTypeAttr : TestAttr<"ParamWithAccessorType"> {
   let parameters = (ins AttrParameter<"std::string", "", "StringRef">:$param);
 }
 
 // DECL-LABEL: class ParamWithAccessorTypeAttr
 // DECL: StringRef getParam()
 // DEF: ParamWithAccessorTypeAttrStorage
 // DEF: ParamWithAccessorTypeAttrStorage(std::string param)

mlir/test/mlir-tblgen/op-result.td

 
 def OpE : NS_Op<"value_attr_as_result_type", [FirstAttrDerivedResultType]> {
   let arguments = (ins I32:$x, F32Attr:$attr);
   let results = (outs AnyTensor:$y);
 }
 
 // CHECK-LABEL: OpE definitions
 // CHECK: void OpE::build(::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &odsState, ::mlir::ValueRange operands, ::llvm::ArrayRef<::mlir::NamedAttribute> attributes)
-// CHECK: odsState.addTypes({attr.getValue().getType()});
+// CHECK: odsState.addTypes({attr.getValue().cast<::mlir::TypedAttr>().getType()});
 
 def OpF : NS_Op<"one_variadic_result_op", []> {
   let results = (outs Variadic<I32>:$x);
 }
 
 // CHECK-LABEL: void OpF::build
 // CHECK-SAME:    ::mlir::TypeRange x
 // CHECK-NOT:     assert

 def OpL3 : NS_Op<"op_with_all_types_constraint",
     [AllTypesMatch<["a", "b"]>]> {
   let arguments = (ins I32Attr:$a);
   let results = (outs AnyType:$b);
 }
 
 // CHECK-LABEL: LogicalResult OpL3::inferReturnTypes
 // CHECK-NOT: }
-// CHECK: ::mlir::Type odsInferredType0 = attributes.get("a").getType();
+// CHECK: ::mlir::Type odsInferredType0 = attributes.get("a").cast<::mlir::TypedAttr>().getType();
 // CHECK: inferredReturnTypes[0] = odsInferredType0;

mlir/tools/mlir-tblgen/AttrOrTypeDefGen.cpp

void DefGen::emitAccessors() {
   for (auto &param : params) {
     Method *m = defCls.addMethod(
         param.getCppAccessorType(), param.getAccessorName(),
         def.genStorageClass() ? Method::Const : Method::ConstDeclaration);
     // Generate accessor definitions only if we also generate the storage
     // class. Otherwise, let the user define the exact accessor definition.
     if (!def.genStorageClass())
       continue;
-    auto scope = m->body().indent().scope("return getImpl()->", ";");
-    if (isa<AttributeSelfTypeParameter>(param))
-      m->body() << formatv("getType().cast<{0}>()", param.getCppType());
-    else
-      m->body() << param.getName();
+    m->body().indent() << "return getImpl()->" << param.getName() << ";";
   }
 }
 
 void DefGen::emitInterfaceMethods() {
   for (auto &traitDef : def.getTraits())
     if (auto *trait = dyn_cast<InterfaceTrait>(&traitDef))
       if (trait->shouldDeclareMethods())
         emitTraitMethods(*trait);

 }
 
 //===----------------------------------------------------------------------===//
 // Storage Class Emission
 
 void DefGen::emitStorageConstructor() {
   Constructor *ctor =
       storageCls->addConstructor<Method::Inline>(getBuilderParams({}));
-  if (auto *attrDef = dyn_cast<AttrDef>(&def)) {
-    // For attributes, a parameter marked with AttributeSelfTypeParameter is
-    // the type initializer that must be passed to the parent constructor.
-    const auto isSelfType = [](const AttrOrTypeParameter &param) {
-      return isa<AttributeSelfTypeParameter>(param);
-    };
-    auto *selfTypeParam = llvm::find_if(params, isSelfType);
-    if (std::count_if(selfTypeParam, params.end(), isSelfType) > 1) {
-      PrintFatalError(def.getLoc(),
-                      "Only one attribute parameter can be marked as "
-                      "AttributeSelfTypeParameter");
-    }
-    // Alternatively, if a type builder was specified, use that instead.
-    std::string attrStorageInit =
-        selfTypeParam == params.end() ? "" : selfTypeParam->getName().str();
-    if (attrDef->getTypeBuilder()) {
-      FmtContext ctx;
-      for (auto &param : params)
-        ctx.addSubst(strfmt("_{0}", param.getName()), param.getName());
-      attrStorageInit = tgfmt(*attrDef->getTypeBuilder(), &ctx);
-    }
-    ctor->addMemberInitializer("::mlir::AttributeStorage",
-                               std::move(attrStorageInit));
-    // Initialize members that aren't the attribute's type.
-    for (auto &param : params)
-      if (selfTypeParam == params.end() || *selfTypeParam != param)
-        ctor->addMemberInitializer(param.getName(), param.getName());
-  } else {
-    for (auto &param : params)
-      ctor->addMemberInitializer(param.getName(), param.getName());
-  }
 }
 
 void DefGen::emitKeyType() {
   std::string keyType("std::tuple<");
   llvm::raw_string_ostream os(keyType);
   llvm::interleaveComma(params, os,
                         [&](auto &param) { os << param.getCppType(); });
   os << '>';
   storageCls->declare<UsingDeclaration>("KeyTy", std::move(os.str()));
 }
 
 void DefGen::emitEquals() {
   Method *eq = storageCls->addConstMethod<Method::Inline>(
       "bool", "operator==", MethodParameter("const KeyTy &", "tblgenKey"));
   auto &body = eq->body().indent();
   auto scope = body.scope("return (", ");");
   const auto eachFn = [&](auto it) {
-    FmtContext ctx({{"_lhs", isa<AttributeSelfTypeParameter>(it.value())
-                                 ? "getType()"
-                                 : it.value().getName()},
+    FmtContext ctx({{"_lhs", it.value().getName()},
                     {"_rhs", strfmt("std::get<{0}>(tblgenKey)", it.index())}});
     body << tgfmt(it.value().getComparator(), &ctx);
   };
   llvm::interleave(llvm::enumerate(params), body, eachFn, ") && (");
 }
 
 void DefGen::emitHashKey() {
   Method *hash = storageCls->addStaticInlineMethod(
void DefGen::emitStorageClass() {
   // Emit the key hash method.
   emitHashKey();
   // Emit the storage constructor. Just declare it if the user wants to define
   // it themself.
   emitConstruct();
   // Emit the storage class members as public, at the very end of the struct.
   storageCls->finalize();
   for (auto &param : params)
-    if (!isa<AttributeSelfTypeParameter>(param))
-      storageCls->declare<Field>(param.getCppType(), param.getName());
 }
 
 //===----------------------------------------------------------------------===//
 // DefGenerator
 //===----------------------------------------------------------------------===//
 
 namespace {
 /// This struct is the base generator used when processing tablegen interfaces.

mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp

private:
   bool lastWasPunctuation = false;
 };
 } // namespace
 
 //===----------------------------------------------------------------------===//
 // ParserGen
 //===----------------------------------------------------------------------===//
 
+/// Generate a special-case "parser" for an attribute's self type parameter. The
+/// self type parameter has special handling in the assembly format in that it
+/// is derived from the optional trailing colon type after the attribute.
+static void genAttrSelfTypeParser(MethodBody &os, const FmtContext &ctx,
+                                  const AttributeSelfTypeParameter &param) {
+  // "Parser" for an attribute self type parameter that checks the
+  // optionally-parsed trailing colon type.
+  //
+  // $0: The C++ storage class of the type parameter.
+  // $1: The self type parameter name.
+  const char *const selfTypeParser = R"(
+if ($_type) {
+  if (auto reqType = $_type.dyn_cast<$0>()) {
+    _result_$1 = reqType;
+  } else {
+    $_parser.emitError($_loc, "invalid kind of type specified");
+    return {};
+  }
+})";
+
+  // If the attribute self type parameter is required, emit code that emits an
+  // error if the trailing type was not parsed.
+  const char *const selfTypeRequired = R"( else {
+  $_parser.emitError($_loc, "expected a trailing type");
+  return {};
+})";
+
+  os << tgfmt(selfTypeParser, &ctx, param.getCppStorageType(), param.getName());
+  if (!param.isOptional())
+    os << tgfmt(selfTypeRequired, &ctx);
+  os << "\n";
+}
+
 void DefFormat::genParser(MethodBody &os) {
   FmtContext ctx;
   ctx.addSubst("_parser", "odsParser");
   ctx.addSubst("_ctxt", "odsParser.getContext()");
   ctx.withBuilder("odsBuilder");
   if (isa<AttrDef>(def))
     ctx.addSubst("_type", "odsType");
   os.indent();
   os << "::mlir::Builder odsBuilder(odsParser.getContext());\n";
 
   // Declare variables to store all of the parameters. Allocated parameters
   // such as `ArrayRef` and `StringRef` must provide a `storageType`. Store
   // FailureOr<T> to defer type construction for parameters that are parsed in
   // a loop (parsers return FailureOr anyways).
   ArrayRef<AttrOrTypeParameter> params = def.getParameters();
   for (const AttrOrTypeParameter &param : params) {
-    if (isa<AttributeSelfTypeParameter>(param))
-      continue;
     os << formatv("::mlir::FailureOr<{0}> _result_{1};\n",
                   param.getCppStorageType(), param.getName());
   }
 
   // Store the initial location of the parser.
   ctx.addSubst("_loc", "odsLoc");
   os << tgfmt("::llvm::SMLoc $_loc = $_parser.getCurrentLocation();\n"
               "(void) $_loc;\n",
               &ctx);
 
   // Generate call to each parameter parser.
   for (FormatElement *el : elements)
     genElementParser(el, ctx, os);
 
   // Emit an assert for each mandatory parameter. Triggering an assert means
   // the generated parser is incorrect (i.e. there is a bug in this code).
   for (const AttrOrTypeParameter &param : params) {
-    if (param.isOptional() || isa<AttributeSelfTypeParameter>(param))
+    if (auto *selfTypeParam = dyn_cast<AttributeSelfTypeParameter>(&param))
+      genAttrSelfTypeParser(os, ctx, *selfTypeParam);
+    if (param.isOptional())
       continue;
     os << formatv("assert(::mlir::succeeded(_result_{0}));\n", param.getName());
   }
 
   // Generate call to the attribute or type builder. Use the checked getter
   // if one was generated.
   if (def.genVerifyDecl()) {
     os << tgfmt("return $_parser.getChecked<$0>($_loc, $_parser.getContext()",
                 &ctx, def.getCppClassName());
   } else {
     os << tgfmt("return $0::get($_parser.getContext()", &ctx,
                 def.getCppClassName());
   }
   for (const AttrOrTypeParameter &param : params) {
     os << ",\n    ";
     std::string paramSelfStr;
     llvm::raw_string_ostream selfOs(paramSelfStr);
     if (param.isOptional()) {
       selfOs << formatv("(_result_{0}.value_or(", param.getName());
       if (Optional<StringRef> defaultValue = param.getDefaultValue())
         selfOs << tgfmt(*defaultValue, &ctx);
       else
         selfOs << param.getCppStorageType() << "()";
       selfOs << "))";
-    } else if (isa<AttributeSelfTypeParameter>(param)) {
-      selfOs << tgfmt("$_type", &ctx);
     } else {
       selfOs << formatv("(*_result_{0})", param.getName());
     }
+    ctx.addSubst(param.getName(), selfOs.str());
     os << param.getCppType() << "("
        << tgfmt(param.getConvertFromStorage(), &ctx.withSelf(selfOs.str()))
        << ")";
   }
   os << ");";
 }
 
 void DefFormat::genElementParser(FormatElement *el, FmtContext &ctx,

mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp

 
 // Populate the format context `ctx` with substitutions of attributes, operands
 // and results.
 static void populateSubstitutions(const OpOrAdaptorHelper &emitHelper,
                                   FmtContext &ctx) {
   // Populate substitutions for attributes.
   auto &op = emitHelper.getOp();
   for (const auto &namedAttr : op.getAttributes())
-    ctx.addSubst(namedAttr.name, emitHelper.getAttr(namedAttr.name).str());
+    ctx.addSubst(namedAttr.name,
+                 emitHelper.getOp().getGetterName(namedAttr.name) + "()");
 
   // Populate substitutions for named operands.
   for (int i = 0, e = op.getNumOperands(); i < e; ++i) {
     auto &value = op.getOperand(i);
     if (!value.name.empty())
       ctx.addSubst(value.name, emitHelper.getOperand(i).str());
   }
 
void OpEmitter::genUseAttrAsResultTypeBuilder() {
   body << "  auto attrName = " << op.getGetterName(namedAttr.name)
        << "AttrName(" << builderOpState
        << ".name);\n"
           "  for (auto attr : attributes) {\n"
           "    if (attr.getName() != attrName) continue;\n";
   if (namedAttr.attr.isTypeAttr()) {
     resultType = "attr.getValue().cast<::mlir::TypeAttr>().getValue()";
   } else {
-    resultType = "attr.getValue().getType()";
+    resultType = "attr.getValue().cast<::mlir::TypedAttr>().getType()";
   }
 
   // Operands
   body << "  " << builderOpState << ".addOperands(operands);\n";
 
   // Attributes
   body << "  " << builderOpState << ".addAttributes(attributes);\n";
 
for (int i = 0, e = op.getNumResults(); i != e; ++i) {
     auto arg = op.getArgToOperandOrAttribute(argIndex);
     if (arg.kind() == Operator::OperandOrAttribute::Kind::Operand) {
       body << "operands[" << arg.operandOrAttributeIndex() << "].getType()";
 
       // If this is an attribute, index into the attribute dictionary.
     } else {
       auto *attr =
           op.getArg(arg.operandOrAttributeIndex()).get<NamedAttribute *>();
-      body << "attributes.get(\"" << attr->name << "\").getType()";
+      body << "attributes.get(\"" << attr->name
+           << "\").cast<::mlir::TypedAttr>().getType()";
     }
     body << ";\n";
   }
 
   // Perform a second pass that handles assigning the inferred types to the
   // results.
   for (int i = 0, e = op.getNumResults(); i != e; ++i) {
     auto types = op.getSameTypeAsResult(i);

mlir/unittests/IR/AttributeTest.cpp

TEST(SparseElementsAttrTest, GetZero) {
 
   auto sparseInt = SparseElementsAttr::get(tensorI32, indices, intValue);
   auto sparseFloat = SparseElementsAttr::get(tensorF32, indices, floatValue);
   auto sparseString =
       SparseElementsAttr::get(tensorString, indices, stringValue);
 
   // Only index (0, 0) contains an element, others are supposed to return
   // the zero/empty value.
-  auto zeroIntValue = sparseInt.getValues<Attribute>()[{1, 1}];
-  EXPECT_EQ(zeroIntValue.cast<IntegerAttr>().getInt(), 0);
+  auto zeroIntValue =
+      sparseInt.getValues<Attribute>()[{1, 1}].cast<IntegerAttr>();
+  EXPECT_EQ(zeroIntValue.getInt(), 0);
   EXPECT_TRUE(zeroIntValue.getType() == intTy);
 
-  auto zeroFloatValue = sparseFloat.getValues<Attribute>()[{1, 1}];
-  EXPECT_EQ(zeroFloatValue.cast<FloatAttr>().getValueAsDouble(), 0.0f);
+  auto zeroFloatValue =
+      sparseFloat.getValues<Attribute>()[{1, 1}].cast<FloatAttr>();
+  EXPECT_EQ(zeroFloatValue.getValueAsDouble(), 0.0f);
   EXPECT_TRUE(zeroFloatValue.getType() == floatTy);
 
-  auto zeroStringValue = sparseString.getValues<Attribute>()[{1, 1}];
-  EXPECT_TRUE(zeroStringValue.cast<StringAttr>().getValue().empty());
+  auto zeroStringValue =
+      sparseString.getValues<Attribute>()[{1, 1}].cast<StringAttr>();
+  EXPECT_TRUE(zeroStringValue.getValue().empty());
   EXPECT_TRUE(zeroStringValue.getType() == stringTy);
 }
 
 } // namespace