[mlir] Populate default attributes on op creation

[jpienaar]

Summary

Default attributes were only handled by ODS accessors generated with the
intention that these behave as if set attributes. This addresses the
long standing TODO to address this inconsistency. Moving the
initialization to construction vs every access. Removing need for
duplicated default attribute population in python bindings.

Switch OpenMP ones to optional attribute with default as the currently
set default values are not legal.

Switched LinAlg generated ones to optional attribute with default as its
quite widely used and unclear where it falls on two different
interpretations.

[rriddle]

LG, did you see any compile time regressions of any kind? Just making sure that our current way of setting the attributes isn't slow for some reason.

[jpienaar]

In D130916#3713426, @rriddle wrote:

LG, did you see any compile time regressions of any kind? Just making sure that our current way of setting the attributes isn't slow for some reason.

I did not no. Most default instances were the simple/empty case though. This should result in no more to less attribute creation unless an accessor is never used (and we could further improve it in future by making the default be cached per op type).

mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td

def CriticalDeclareOp : OpenMP_Op<"critical.declare", [Symbol]> {
 
   let description = [{
     Declares a named critical section.
 
     The name can be used in critical constructs in the dialect.
   }];
 
   let arguments = (ins SymbolNameAttr:$sym_name,
-                       DefaultValuedAttr<I64Attr, "0">:$hint_val);
+                       DefaultValuedOptionalAttr<I64Attr, "0">:$hint_val);
 
   let assemblyFormat = [{
     $sym_name oilist(`hint` `(` custom<SynchronizationHint>($hint_val) `)`)
     attr-dict
   }];
   let hasVerifier = 1;
 }
 
let description = [{
     optimization.
 
     `memory_order` indicates the memory ordering behavior of the construct. It
     can be one of `seq_cst`, `acquire` or `relaxed`.
   }];
 
   let arguments = (ins OpenMP_PointerLikeType:$x,
                        OpenMP_PointerLikeType:$v,
-                       DefaultValuedAttr<I64Attr, "0">:$hint_val,
+                       DefaultValuedOptionalAttr<I64Attr, "0">:$hint_val,
                        OptionalAttr<MemoryOrderKindAttr>:$memory_order_val);
   let assemblyFormat = [{
     $v `=` $x 
     oilist( `memory_order` `(` custom<ClauseAttr>($memory_order_val) `)`
           | `hint` `(` custom<SynchronizationHint>($hint_val) `)`)
     `:` type($x) attr-dict
   }];
   let hasVerifier = 1;
let description = [{
     optimization.
 
     `memory_order` indicates the memory ordering behavior of the construct. It
     can be one of `seq_cst`, `release` or `relaxed`.
   }];
 
   let arguments = (ins OpenMP_PointerLikeType:$address,
                        AnyType:$value,
-                       DefaultValuedAttr<I64Attr, "0">:$hint_val,
+                       DefaultValuedOptionalAttr<I64Attr, "0">:$hint_val,
                        OptionalAttr<MemoryOrderKindAttr>:$memory_order_val);
   let assemblyFormat = [{
     $address `=` $value
     oilist( `hint` `(` custom<SynchronizationHint>($hint_val) `)`
           | `memory_order` `(` custom<ClauseAttr>($memory_order_val) `)`)
     `:` type($address) `,` type($value)
     attr-dict
   }];
let description = [{
     the core update operation is directly translated like regular operations by
     the host dialect. The front-end must handle semantic checks for allowed
     operations.
   }];
 
   let arguments = (ins Arg<OpenMP_PointerLikeType,
                            "Address of variable to be updated",
                            [MemRead, MemWrite]>:$x,
-                       DefaultValuedAttr<I64Attr, "0">:$hint_val,
+                       DefaultValuedOptionalAttr<I64Attr, "0">:$hint_val,
                        OptionalAttr<MemoryOrderKindAttr>:$memory_order_val);
   let regions = (region SizedRegion<1>:$region);
   let assemblyFormat = [{
     oilist( `memory_order` `(` custom<ClauseAttr>($memory_order_val) `)`
           | `hint` `(` custom<SynchronizationHint>($hint_val) `)`)
     $x `:` type($x) $region attr-dict
   }];
   let hasVerifier = 1;
```
         omp.atomic.read ...
         omp.atomic.write ...
         omp.terminator
       }
     ```
 
   }];
 
-  let arguments = (ins DefaultValuedAttr<I64Attr, "0">:$hint_val,
+  let arguments = (ins DefaultValuedOptionalAttr<I64Attr, "0">:$hint_val,
                        OptionalAttr<MemoryOrderKindAttr>:$memory_order_val);
   let regions = (region SizedRegion<1>:$region);
   let assemblyFormat = [{
     oilist(`memory_order` `(` custom<ClauseAttr>($memory_order_val) `)`
           |`hint` `(` custom<SynchronizationHint>($hint_val) `)`)
     $region attr-dict
   }];
   let hasRegionVerifier = 1;

mlir/lib/IR/Operation.cpp

size_t byteSize =
       totalSizeToAlloc<detail::OperandStorage, BlockOperand, Region, OpOperand>(
           needsOperandStorage ? 1 : 0, numSuccessors, numRegions, numOperands);
   size_t prefixByteSize = llvm::alignTo(
       Operation::prefixAllocSize(numTrailingResults, numInlineResults),
       alignof(Operation));
   char *mallocMem = reinterpret_cast<char *>(malloc(byteSize + prefixByteSize));
   void *rawMem = mallocMem + prefixByteSize;
 
+  // Populate default attributes.
+  if (Optional<RegisteredOperationName> info = name.getRegisteredInfo())
+    info->populateDefaultAttrs(attributes);
+
   // Create the new Operation.
   Operation *op = ::new (rawMem) Operation(
       location, name, numResults, numSuccessors, numRegions,
       attributes.getDictionary(location.getContext()), needsOperandStorage);
 
   assert((numSuccessors == 0 || op->mightHaveTrait<OpTrait::IsTerminator>()) &&
          "unexpected successors in a non-terminator operation");
 

mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir

 //       CHECK:    llvm.insertvalue {{.*}}[3, 1] : !llvm.struct<(ptr<f32>, ptr<f32>, i64, array<3 x i64>, array<3 x i64>)>
 func.func @transpose(%arg0: memref<?x?x?xf32, offset: ?, strides: [?, ?, 1]>) {
   %0 = memref.transpose %arg0 (i, j, k) -> (k, i, j) : memref<?x?x?xf32, offset: ?, strides: [?, ?, 1]> to memref<?x?x?xf32, affine_map<(d0, d1, d2)[s0, s1, s2] -> (d2 * s1 + s0 + d0 * s2 + d1)>>
   return
 }
 
 // -----
 
-// CHECK:   llvm.mlir.global external @gv0() : !llvm.array<2 x f32> {
+// CHECK:   llvm.mlir.global external @gv0() {addr_space = 0 : i32} : !llvm.array<2 x f32> {
 // CHECK-NEXT:     %0 = llvm.mlir.undef : !llvm.array<2 x f32>
 // CHECK-NEXT:     llvm.return %0 : !llvm.array<2 x f32>
 // CHECK-NEXT:   }
 memref.global @gv0 : memref<2xf32> = uninitialized
 
-// CHECK: llvm.mlir.global private @gv1() : !llvm.array<2 x f32>
+// CHECK: llvm.mlir.global private @gv1() {addr_space = 0 : i32} : !llvm.array<2 x f32>
 memref.global "private" @gv1 : memref<2xf32>
 
-// CHECK: llvm.mlir.global external @gv2(dense<{{\[\[}}0.000000e+00, 1.000000e+00, 2.000000e+00], [3.000000e+00, 4.000000e+00, 5.000000e+00]]> : tensor<2x3xf32>) : !llvm.array<2 x array<3 x f32>>
+// CHECK: llvm.mlir.global external @gv2(dense<{{\[\[}}0.000000e+00, 1.000000e+00, 2.000000e+00], [3.000000e+00, 4.000000e+00, 5.000000e+00]]> : tensor<2x3xf32>) {addr_space = 0 : i32} : !llvm.array<2 x array<3 x f32>>
 memref.global @gv2 : memref<2x3xf32> = dense<[[0.0, 1.0, 2.0], [3.0, 4.0, 5.0]]>
 
 // Test 1D memref.
 // CHECK-LABEL: func @get_gv0_memref
 func.func @get_gv0_memref() {
   %0 = memref.get_global @gv0 : memref<2xf32>
   // CHECK: %[[DIM:.*]] = llvm.mlir.constant(2 : index) : i64
   // CHECK: %[[STRIDE:.*]] = llvm.mlir.constant(1 : index) : i64
func.func @get_gv2_memref() {
   // CHECK: llvm.insertvalue %[[DIM1]], {{.*}}[4, 0] : !llvm.struct<(ptr<f32>, ptr<f32>, i64, array<2 x i64>, array<2 x i64>)>
   // CHECK: llvm.insertvalue %[[STRIDE1]], {{.*}}[4, 1] : !llvm.struct<(ptr<f32>, ptr<f32>, i64, array<2 x i64>, array<2 x i64>)>
 
   %0 = memref.get_global @gv2 : memref<2x3xf32>
   return
 }
 
 // Test scalar memref.
-// CHECK: llvm.mlir.global external @gv3(1.000000e+00 : f32) : f32
+// CHECK: llvm.mlir.global external @gv3(1.000000e+00 : f32) {addr_space = 0 : i32} : f32
 memref.global @gv3 : memref<f32> = dense<1.0>
 
 // CHECK-LABEL: func @get_gv3_memref
 func.func @get_gv3_memref() {
   // CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @gv3 : !llvm.ptr<f32>
   // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : index) : i64
   // CHECK: %[[GEP:.*]] = llvm.getelementptr %[[ADDR]][%[[ZERO]]] : (!llvm.ptr<f32>, i64) -> !llvm.ptr<f32>
   // CHECK: %[[DEADBEEF:.*]] = llvm.mlir.constant(3735928559 : index) : i64
   // CHECK: %[[DEADBEEFPTR:.*]] = llvm.inttoptr %[[DEADBEEF]] : i64 to !llvm.ptr<f32>
   // CHECK: llvm.mlir.undef : !llvm.struct<(ptr<f32>, ptr<f32>, i64)>
   // CHECK: llvm.insertvalue %[[DEADBEEFPTR]], {{.*}}[0] : !llvm.struct<(ptr<f32>, ptr<f32>, i64)>
   // CHECK: llvm.insertvalue %[[GEP]], {{.*}}[1] : !llvm.struct<(ptr<f32>, ptr<f32>, i64)>
   // CHECK: %[[OFFSET:.*]] = llvm.mlir.constant(0 : index) : i64
   // CHECK: llvm.insertvalue %[[OFFSET]], {{.*}}[2] : !llvm.struct<(ptr<f32>, ptr<f32>, i64)>
   %0 = memref.get_global @gv3 : memref<f32>
   return
 }
 
 // Test scalar memref with an alignment.
-// CHECK: llvm.mlir.global private @gv4(1.000000e+00 : f32) {alignment = 64 : i64} : f32
+// CHECK: llvm.mlir.global private @gv4(1.000000e+00 : f32) {addr_space = 0 : i32, alignment = 64 : i64} : f32
 memref.global "private" @gv4 : memref<f32> = dense<1.0> {alignment = 64}
 
 // -----
 
 func.func @collapse_shape_static(%arg0: memref<1x3x4x1x5xf32>) -> memref<3x4x5xf32> {
   %0 = memref.collapse_shape %arg0 [[0, 1], [2], [3, 4]] :
     memref<1x3x4x1x5xf32> into memref<3x4x5xf32>
   return %0 : memref<3x4x5xf32>

mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir

 // RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s
 
 module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, #spv.resource_limits<max_compute_workgroup_invocations = 128, max_compute_workgroup_size = [128, 128, 64]>>} {
 
-  //       CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() : !llvm.struct<(array<6 x i32>)>
+  //       CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() {addr_space = 0 : i32} : !llvm.struct<(array<6 x i32>)>
   //       CHECK: llvm.func @__spv__foo_bar()
 
   //       CHECK: spv.module @__spv__foo
   //       CHECK:   spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
   //       CHECK:   spv.func @__spv__foo_bar
 
   //       CHECK:   spv.EntryPoint "GLCompute" @__spv__foo_bar
   //       CHECK:   spv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1

mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir

spv.func @access_chain_array(%arg0 : i32) "None" {
   spv.Return
 }
 
 //===----------------------------------------------------------------------===//
 // spv.GlobalVariable and spv.mlir.addressof
 //===----------------------------------------------------------------------===//
 
 spv.module Logical GLSL450 {
-  // CHECK: llvm.mlir.global external constant @var() : f32
+  // CHECK: llvm.mlir.global external constant @var() {addr_space = 0 : i32} : f32
   spv.GlobalVariable @var : !spv.ptr<f32, Input>
 }
 
 spv.module Logical GLSL450 {
-  //       CHECK: llvm.mlir.global private @struct() : !llvm.struct<packed (f32, array<10 x f32>)>
+  //       CHECK: llvm.mlir.global private @struct() {addr_space = 0 : i32} : !llvm.struct<packed (f32, array<10 x f32>)>
   // CHECK-LABEL: @func
   //       CHECK:   llvm.mlir.addressof @struct : !llvm.ptr<struct<packed (f32, array<10 x f32>)>>
   spv.GlobalVariable @struct : !spv.ptr<!spv.struct<(f32, !spv.array<10xf32>)>, Private>
   spv.func @func() "None" {
     %0 = spv.mlir.addressof @struct : !spv.ptr<!spv.struct<(f32, !spv.array<10xf32>)>, Private>
     spv.Return
   }
 }
 
 spv.module Logical GLSL450 {
-  //       CHECK: llvm.mlir.global external @bar_descriptor_set0_binding0() : i32
+  //       CHECK: llvm.mlir.global external @bar_descriptor_set0_binding0() {addr_space = 0 : i32} : i32
   // CHECK-LABEL: @foo
   //       CHECK:   llvm.mlir.addressof @bar_descriptor_set0_binding0 : !llvm.ptr<i32>
   spv.GlobalVariable @bar bind(0, 0) : !spv.ptr<i32, StorageBuffer>
   spv.func @foo() "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<i32, StorageBuffer>
     spv.Return
   }
 }
 
 spv.module @name Logical GLSL450 {
-  //       CHECK: llvm.mlir.global external @name_bar_descriptor_set0_binding0() : i32
+  //       CHECK: llvm.mlir.global external @name_bar_descriptor_set0_binding0() {addr_space = 0 : i32} : i32
   // CHECK-LABEL: @foo
   //       CHECK:   llvm.mlir.addressof @name_bar_descriptor_set0_binding0 : !llvm.ptr<i32>
   spv.GlobalVariable @bar bind(0, 0) : !spv.ptr<i32, StorageBuffer>
   spv.func @foo() "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<i32, StorageBuffer>
     spv.Return
   }
 }
 
 spv.module Logical GLSL450 {
-  // CHECK: llvm.mlir.global external @bar() {location = 1 : i32} : i32
+  // CHECK: llvm.mlir.global external @bar() {addr_space = 0 : i32, location = 1 : i32} : i32
   // CHECK-LABEL: @foo
   spv.GlobalVariable @bar {location = 1 : i32} : !spv.ptr<i32, Output>
   spv.func @foo() "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<i32, Output>
     spv.Return
   }
 }
 
 spv.module Logical GLSL450 {
-  // CHECK: llvm.mlir.global external constant @bar() {location = 3 : i32} : f32
+  // CHECK: llvm.mlir.global external constant @bar() {addr_space = 0 : i32, location = 3 : i32} : f32
   // CHECK-LABEL: @foo
   spv.GlobalVariable @bar {descriptor_set = 0 : i32, location = 3 : i32} : !spv.ptr<f32, UniformConstant>
   spv.func @foo() "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<f32, UniformConstant>
     spv.Return
   }
 }
 

mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir

spv.func @vector_shuffle_different_size(%vector1: vector<3xf32>, %vector2: vector<2xf32>) -> vector<3xf32> "None" {
   spv.ReturnValue %0: vector<3xf32>
 }
 
 //===----------------------------------------------------------------------===//
 // spv.EntryPoint and spv.ExecutionMode
 //===----------------------------------------------------------------------===//
 
 //      CHECK: module {
-// CHECK-NEXT:   llvm.mlir.global external constant @{{.*}}() : !llvm.struct<(i32)> {
+// CHECK-NEXT:   llvm.mlir.global external constant @{{.*}}() {addr_space = 0 : i32} : !llvm.struct<(i32)> {
 // CHECK-NEXT:     %[[UNDEF:.*]] = llvm.mlir.undef : !llvm.struct<(i32)>
 // CHECK-NEXT:     %[[VAL:.*]] = llvm.mlir.constant(31 : i32) : i32
 // CHECK-NEXT:     %[[RET:.*]] = llvm.insertvalue %[[VAL]], %[[UNDEF]][0 : i32] : !llvm.struct<(i32)>
 // CHECK-NEXT:     llvm.return %[[RET]] : !llvm.struct<(i32)>
 // CHECK-NEXT:   }
 // CHECK-NEXT:   llvm.func @empty
 // CHECK-NEXT:     llvm.return
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 spv.module Logical OpenCL {
   spv.func @empty() "None" {
     spv.Return
   }
   spv.EntryPoint "Kernel" @empty
   spv.ExecutionMode @empty "ContractionOff"
 }
 
 //      CHECK: module {
-// CHECK-NEXT:   llvm.mlir.global external constant @{{.*}}() : !llvm.struct<(i32, array<3 x i32>)> {
+// CHECK-NEXT:   llvm.mlir.global external constant @{{.*}}() {addr_space = 0 : i32} : !llvm.struct<(i32, array<3 x i32>)> {
 // CHECK-NEXT:     %[[UNDEF:.*]] = llvm.mlir.undef : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:     %[[EM:.*]] = llvm.mlir.constant(18 : i32) : i32
 // CHECK-NEXT:     %[[T0:.*]] = llvm.insertvalue %[[EM]], %[[UNDEF]][0 : i32] : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:     %[[C0:.*]] = llvm.mlir.constant(32 : i32) : i32
 // CHECK-NEXT:     %[[T1:.*]] = llvm.insertvalue %[[C0]], %[[T0]][1 : i32, 0 : i32] : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:     %[[C1:.*]] = llvm.mlir.constant(1 : i32) : i32
 // CHECK-NEXT:     %[[T2:.*]] = llvm.insertvalue %[[C1]], %[[T1]][1 : i32, 1 : i32] : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:     %[[C2:.*]] = llvm.mlir.constant(1 : i32) : i32
 // CHECK-NEXT:     %[[RET:.*]] = llvm.insertvalue %[[C2]], %[[T2]][1 : i32, 2 : i32] : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:     llvm.return %[[RET]] : !llvm.struct<(i32, array<3 x i32>)>
 // CHECK-NEXT:   }
-// CHECK-NEXT:   llvm.mlir.global external constant @{{.*}}() : !llvm.struct<(i32)> {
+// CHECK-NEXT:   llvm.mlir.global external constant @{{.*}}() {addr_space = 0 : i32} : !llvm.struct<(i32)> {
 //      CHECK:   llvm.func @bar
 // CHECK-NEXT:     llvm.return
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 spv.module Logical OpenCL {
   spv.func @bar() "None" {
     spv.Return
   }

mlir/test/Dialect/GPU/outlining.mlir

 
 // CHECK: gpu.module @function_call_kernel {
 // CHECK:   gpu.func @function_call_kernel()
 // CHECK:     call @device_function() : () -> ()
 // CHECK:     call @device_function() : () -> ()
 // CHECK:     llvm.mlir.addressof @global : !llvm.ptr<i64>
 // CHECK:     gpu.return
 //
-// CHECK:   llvm.mlir.global internal @global(42 : i64) : i64
+// CHECK:   llvm.mlir.global internal @global(42 : i64) {addr_space = 0 : i32} : i64
 //
 // CHECK:   func @device_function()
 // CHECK:   func @recursive_device_function()
 // CHECK-NOT:   func @device_function

mlir/test/Dialect/LLVMIR/global.mlir

 // RUN: mlir-opt -split-input-file -verify-diagnostics %s | FileCheck %s
 
 // CHECK: llvm.mlir.global external @default_external
 llvm.mlir.global @default_external() : i64
 
 // CHECK: llvm.mlir.global external constant @default_external_constant
 llvm.mlir.global constant @default_external_constant(42) : i64
 
-// CHECK: llvm.mlir.global internal @global(42 : i64) : i64
+// CHECK: llvm.mlir.global internal @global(42 : i64) {addr_space = 0 : i32} : i64
 llvm.mlir.global internal @global(42 : i64) : i64
 
-// CHECK: llvm.mlir.global private @aligned_global(42 : i64) {aligned = 64 : i64} : i64
+// CHECK: llvm.mlir.global private @aligned_global(42 : i64) {addr_space = 0 : i32, aligned = 64 : i64} : i64
 llvm.mlir.global private @aligned_global(42 : i64) {aligned = 64} : i64
 
-// CHECK: llvm.mlir.global private constant @aligned_global_const(42 : i64) {aligned = 32 : i64} : i64
+// CHECK: llvm.mlir.global private constant @aligned_global_const(42 : i64) {addr_space = 0 : i32, aligned = 32 : i64} : i64
 llvm.mlir.global private constant @aligned_global_const(42 : i64) {aligned = 32} : i64
 
-// CHECK: llvm.mlir.global internal constant @constant(3.700000e+01 : f64) : f32
+// CHECK: llvm.mlir.global internal constant @constant(3.700000e+01 : f64) {addr_space = 0 : i32} : f32
 llvm.mlir.global internal constant @constant(37.0) : f32
 
 // CHECK: llvm.mlir.global internal constant @".string"("foobar")
 llvm.mlir.global internal constant @".string"("foobar") : !llvm.array<6 x i8>
 
 // CHECK: llvm.mlir.global internal @string_notype("1234567")
 llvm.mlir.global internal @string_notype("1234567")
 
 // CHECK: llvm.mlir.global internal @global_undef()
 llvm.mlir.global internal @global_undef() : i64
 
-// CHECK: llvm.mlir.global internal @global_mega_initializer() : i64 {
+// CHECK: llvm.mlir.global internal @global_mega_initializer() {addr_space = 0 : i32} : i64 {
 // CHECK-NEXT:  %[[c:[0-9]+]] = llvm.mlir.constant(42 : i64) : i64
 // CHECK-NEXT:  llvm.return %[[c]] : i64
 // CHECK-NEXT: }
 llvm.mlir.global internal @global_mega_initializer() : i64 {
   %c = llvm.mlir.constant(42 : i64) : i64
   llvm.return %c : i64
 }
 

 // CHECK: llvm.mlir.global appending
 llvm.mlir.global appending @appending() : !llvm.array<2 x i64>
 // CHECK: llvm.mlir.global extern_weak
 llvm.mlir.global extern_weak @extern_weak() : i64
 // CHECK: llvm.mlir.global linkonce_odr
 llvm.mlir.global linkonce_odr @linkonce_odr() : i64
 // CHECK: llvm.mlir.global weak_odr
 llvm.mlir.global weak_odr @weak_odr() : i64
-// CHECK: llvm.mlir.global external @has_thr_local(42 : i64) {thr_local} : i64
+// CHECK: llvm.mlir.global external @has_thr_local(42 : i64) {addr_space = 0 : i32, thr_local} : i64
 llvm.mlir.global external @has_thr_local(42 : i64) {thr_local} : i64
-// CHECK: llvm.mlir.global external @has_dso_local(42 : i64) {dso_local} : i64
+// CHECK: llvm.mlir.global external @has_dso_local(42 : i64) {addr_space = 0 : i32, dso_local} : i64
 llvm.mlir.global external @has_dso_local(42 : i64) {dso_local} : i64
 // CHECK: llvm.mlir.global external @has_addr_space(32 : i64) {addr_space = 3 : i32} : i64
 llvm.mlir.global external @has_addr_space(32 : i64) {addr_space = 3: i32} : i64
 
 // CHECK-LABEL: references
 func.func @references() {
   // CHECK: llvm.mlir.addressof @global : !llvm.ptr<i64>
   %0 = llvm.mlir.addressof @global : !llvm.ptr<i64>
 
   // CHECK: llvm.mlir.addressof @".string" : !llvm.ptr<array<6 x i8>>
   %1 = llvm.mlir.addressof @".string" : !llvm.ptr<array<6 x i8>>
 
   // CHECK: llvm.mlir.addressof @global : !llvm.ptr
   %2 = llvm.mlir.addressof @global : !llvm.ptr
 
   // CHECK: llvm.mlir.addressof @has_addr_space : !llvm.ptr<3>
   %3 = llvm.mlir.addressof @has_addr_space : !llvm.ptr<3>
 
   llvm.return
 }
 
-// CHECK: llvm.mlir.global private local_unnamed_addr constant @local(42 : i64) : i64
+// CHECK: llvm.mlir.global private local_unnamed_addr constant @local(42 : i64) {addr_space = 0 : i32} : i64
 llvm.mlir.global private local_unnamed_addr constant @local(42 : i64) : i64
 
-// CHECK: llvm.mlir.global private unnamed_addr constant @foo(42 : i64) : i64
+// CHECK: llvm.mlir.global private unnamed_addr constant @foo(42 : i64) {addr_space = 0 : i32} : i64
 llvm.mlir.global private unnamed_addr constant @foo(42 : i64) : i64
 
-// CHECK: llvm.mlir.global internal constant @sectionvar("teststring")  {section = ".mysection"}
+// CHECK: llvm.mlir.global internal constant @sectionvar("teststring") {addr_space = 0 : i32, section = ".mysection"}
 llvm.mlir.global internal constant @sectionvar("teststring")  {section = ".mysection"}: !llvm.array<10 x i8>
 
 // -----
 
 // expected-error @+1 {{op requires attribute 'sym_name'}}
 "llvm.mlir.global"() ({}) {linkage = "private", type = i64, constant, global_type = i64, value = 42 : i64} : () -> ()
 
 // -----

mlir/test/Dialect/Linalg/vectorize-convolution.mlir

 //      CHECK:  %[[V_OUTPUT_0:.+]] = vector.extract_strided_slice %[[V_OUTPUT_R]]
 // CHECK-SAME:     {offsets = [0, 0, 0], sizes = [4, 1, 8], strides = [1, 1, 1]} : vector<4x2x8xf32> to vector<4x1x8xf32>
 //      CHECK:  %[[V_OUTPUT_1:.+]] = vector.extract_strided_slice %[[V_OUTPUT_R]]
 // CHECK-SAME:     {offsets = [0, 1, 0], sizes = [4, 1, 8], strides = [1, 1, 1]} : vector<4x2x8xf32> to vector<4x1x8xf32>
 
 /// w == 0, kw == 0
 //      CHECK:   %[[CONTRACT_0:.+]] = vector.contract {
 // CHECK-SAME:       indexing_maps = [#[[INPUT_MAP]], #[[FILTER_MAP]], #[[OUTPUT_MAP]]],
-// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]}
+// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]
 // CHECK-SAME:     %[[V_INPUT_0]], %[[V_FILTER]], %[[V_OUTPUT_0]]
 // CHECK-SAME:     : vector<4x1x3xf32>, vector<3x8xf32> into vector<4x1x8xf32>
 
 /// w == 1, kw == 0
 //      CHECK:   %[[CONTRACT_1:.+]] = vector.contract {
 // CHECK-SAME:       indexing_maps = [#[[INPUT_MAP]], #[[FILTER_MAP]], #[[OUTPUT_MAP]]],
-// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]}
+// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]
 // CHECK-SAME:     %[[V_INPUT_1]], %[[V_FILTER]], %[[V_OUTPUT_1]]
 // CHECK-SAME:     : vector<4x1x3xf32>, vector<3x8xf32> into vector<4x1x8xf32>
 
 /// w == 0, kw == 0
 //      CHECK:   %[[RES_0:.+]] = vector.insert_strided_slice %[[CONTRACT_0]], %[[V_OUTPUT_R]]
 // CHECK-SAME:     {offsets = [0, 0, 0], strides = [1, 1, 1]} : vector<4x1x8xf32> into vector<4x2x8xf32>
 /// w == 1, kw == 0
 //      CHECK:   %[[RES_1:.+]] = vector.insert_strided_slice %[[CONTRACT_1]], %[[RES_0]]

 //      CHECK:  %[[V_OUTPUT_0:.+]] = vector.extract_strided_slice %[[V_OUTPUT_R]]
 // CHECK-SAME:     {offsets = [0, 0, 0], sizes = [4, 1, 8], strides = [1, 1, 1]} : vector<4x2x8xf32> to vector<4x1x8xf32>
 //      CHECK:  %[[V_OUTPUT_1:.+]] = vector.extract_strided_slice %[[V_OUTPUT_R]]
 // CHECK-SAME:     {offsets = [0, 1, 0], sizes = [4, 1, 8], strides = [1, 1, 1]} : vector<4x2x8xf32> to vector<4x1x8xf32>
 
 /// w == 0, kw == 0
 //      CHECK:   %[[CONTRACT_0:.+]] = vector.contract {
 // CHECK-SAME:       indexing_maps = [#[[INPUT_MAP]], #[[FILTER_MAP]], #[[OUTPUT_MAP]]],
-// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]}
+// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]
 // CHECK-SAME:     %[[V_INPUT_0]], %[[V_FILTER_0]], %[[V_OUTPUT_0]]
 // CHECK-SAME:     : vector<4x1x3xf32>, vector<3x8xf32> into vector<4x1x8xf32>
 /// w == 1, kw == 0
 //      CHECK:   %[[CONTRACT_1:.+]] = vector.contract {
 // CHECK-SAME:       indexing_maps = [#[[INPUT_MAP]], #[[FILTER_MAP]], #[[OUTPUT_MAP]]],
-// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]}
+// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]
 // CHECK-SAME:     %[[V_INPUT_1]], %[[V_FILTER_0]], %[[V_OUTPUT_1]]
 // CHECK-SAME:     : vector<4x1x3xf32>, vector<3x8xf32> into vector<4x1x8xf32>
 /// w == 1, kw == 1
 //      CHECK:   %[[CONTRACT_2:.+]] = vector.contract {
 // CHECK-SAME:       indexing_maps = [#[[INPUT_MAP]], #[[FILTER_MAP]], #[[OUTPUT_MAP]]],
-// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]}
+// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]
 // CHECK-SAME:     %[[V_INPUT_2]], %[[V_FILTER_1]], %[[CONTRACT_0]]
 // CHECK-SAME:     : vector<4x1x3xf32>, vector<3x8xf32> into vector<4x1x8xf32>
 /// w == 1, kw == 1
 //      CHECK:   %[[CONTRACT_3:.+]] = vector.contract {
 // CHECK-SAME:       indexing_maps = [#[[INPUT_MAP]], #[[FILTER_MAP]], #[[OUTPUT_MAP]]],
-// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]}
+// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]
 // CHECK-SAME:     %[[V_INPUT_3]], %[[V_FILTER_1]], %[[CONTRACT_1]]
 // CHECK-SAME:     : vector<4x1x3xf32>, vector<3x8xf32> into vector<4x1x8xf32>
 
 /// w == 0, kw == 0
 //      CHECK:   %[[RES_0:.+]] = vector.insert_strided_slice %[[CONTRACT_2]], %[[V_OUTPUT_R]]
 // CHECK-SAME:     {offsets = [0, 0, 0], strides = [1, 1, 1]} : vector<4x1x8xf32> into vector<4x2x8xf32>
 /// w == 1, kw == 0
 //      CHECK:   %[[RES_1:.+]] = vector.insert_strided_slice %[[CONTRACT_3]], %[[RES_0]]

 // CHECK-SAME:     {offsets = [0, 2, 0], sizes = [4, 2, 3], strides = [1, 1, 1]} : vector<4x4x3xf32> to vector<4x2x3xf32>
 
 //      CHECK:  %[[V_FILTER_0:.+]] = vector.extract %[[V_FILTER_R]][0] : vector<2x3x8xf32>
 //      CHECK:  %[[V_FILTER_1:.+]] = vector.extract %[[V_FILTER_R]][1] : vector<2x3x8xf32>
 
 /// w == 0, kw == 0
 //      CHECK:   %[[CONTRACT_0:.+]] = vector.contract {
 // CHECK-SAME:       indexing_maps = [#[[INPUT_MAP]], #[[FILTER_MAP]], #[[OUTPUT_MAP]]],
-// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]}
+// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]
 // CHECK-SAME:     %[[V_INPUT_0]], %[[V_FILTER_0]], %[[V_OUTPUT_R]]
 // CHECK-SAME:     : vector<4x2x3xf32>, vector<3x8xf32> into vector<4x2x8xf32>
 /// w == 0, kw == 1
 //      CHECK:   %[[CONTRACT_1:.+]] = vector.contract {
 // CHECK-SAME:       indexing_maps = [#[[INPUT_MAP]], #[[FILTER_MAP]], #[[OUTPUT_MAP]]],
-// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]}
+// CHECK-SAME:       iterator_types = ["parallel", "parallel", "parallel", "reduction"]
 // CHECK-SAME:     %[[V_INPUT_1]], %[[V_FILTER_1]], %[[CONTRACT_0]]
 // CHECK-SAME:     : vector<4x2x3xf32>, vector<3x8xf32> into vector<4x2x8xf32>
 
 // Write the result back in one shot.
 //      CHECK:   vector.transfer_write %[[CONTRACT_1]], %[[OUTPUT]][%[[C0]], %[[C0]], %[[C0]]]
 
 // -----
 

mlir/test/Target/LLVMIR/Import/basic.ll

 ; RUN: mlir-translate -import-llvm %s | FileCheck %s
 
 %struct.t = type {}
 %struct.s = type { %struct.t, i64 }
 
-; CHECK: llvm.mlir.global external @g1() {alignment = 8 : i64} : !llvm.struct<"struct.s", (struct<"struct.t", ()>, i64)>
+; CHECK: llvm.mlir.global external @g1() {addr_space = 0 : i32, alignment = 8 : i64} : !llvm.struct<"struct.s", (struct<"struct.t", ()>, i64)>
 @g1 = external global %struct.s, align 8
-; CHECK: llvm.mlir.global external @g2() {alignment = 8 : i64} : f64
+; CHECK: llvm.mlir.global external @g2() {addr_space = 0 : i32, alignment = 8 : i64} : f64
 @g2 = external global double, align 8
 ; CHECK: llvm.mlir.global internal @g3("string")
 @g3 = internal global [6 x i8] c"string"
 
-; CHECK: llvm.mlir.global external @g5() : vector<8xi32>
+; CHECK: llvm.mlir.global external @g5() {addr_space = 0 : i32} : vector<8xi32>
 @g5 = external global <8 x i32>
 
-; CHECK: llvm.mlir.global private @alig32(42 : i64) {alignment = 32 : i64, dso_local} : i64
+; CHECK: llvm.mlir.global private @alig32(42 : i64) {addr_space = 0 : i32, alignment = 32 : i64, dso_local} : i64
 @alig32 = private global i64 42, align 32
 
-; CHECK: llvm.mlir.global private @alig64(42 : i64) {alignment = 64 : i64, dso_local} : i64
+; CHECK: llvm.mlir.global private @alig64(42 : i64) {addr_space = 0 : i32, alignment = 64 : i64, dso_local} : i64
 @alig64 = private global i64 42, align 64
 
 @g4 = external global i32, align 8
-; CHECK: llvm.mlir.global internal constant @int_gep() {dso_local} : !llvm.ptr<i32> {
+; CHECK: llvm.mlir.global internal constant @int_gep() {addr_space = 0 : i32, dso_local} : !llvm.ptr<i32> {
 ; CHECK-DAG:   %[[addr:[0-9]+]] = llvm.mlir.addressof @g4 : !llvm.ptr<i32>
 ; CHECK-DAG:   %[[c2:[0-9]+]] = llvm.mlir.constant(2 : i32) : i32
 ; CHECK-NEXT:  %[[gepinit:[0-9]+]] = llvm.getelementptr %[[addr]][%[[c2]]] : (!llvm.ptr<i32>, i32) -> !llvm.ptr<i32>
 ; CHECK-NEXT:  llvm.return %[[gepinit]] : !llvm.ptr<i32>
 ; CHECK-NEXT: }
 @int_gep = internal constant i32* getelementptr (i32, i32* @g4, i32 2)
 
 ;
 ; dso_local attribute
 ;
 
-; CHECK: llvm.mlir.global external @dso_local_var() {dso_local} : !llvm.struct<"struct.s", (struct<"struct.t", ()>, i64)>
+; CHECK: llvm.mlir.global external @dso_local_var() {addr_space = 0 : i32, dso_local} : !llvm.struct<"struct.s", (struct<"struct.t", ()>, i64)>
 @dso_local_var = external dso_local global %struct.s
 
 ;
 ; thread_local attribute
 ;
 
-; CHECK: llvm.mlir.global external thread_local @thread_local_var() : !llvm.struct<"struct.s", (struct<"struct.t", ()>, i64)>
+; CHECK: llvm.mlir.global external thread_local @thread_local_var() {addr_space = 0 : i32} : !llvm.struct<"struct.s", (struct<"struct.t", ()>, i64)>
 @thread_local_var = external thread_local global %struct.s
 
 ;
 ; addr_space attribute
 ;
 
 ; CHECK: llvm.mlir.global external @addr_space_var(0 : i32) {addr_space = 6 : i32} : i32
 @addr_space_var = addrspace(6) global i32 0
 
 ;
 ; Linkage attribute.
 ;
 
-; CHECK: llvm.mlir.global private @private(42 : i32) {dso_local} : i32
+; CHECK: llvm.mlir.global private @private(42 : i32) {addr_space = 0 : i32, dso_local} : i32
 @private = private global i32 42
-; CHECK: llvm.mlir.global internal @internal(42 : i32) {dso_local} : i32
+; CHECK: llvm.mlir.global internal @internal(42 : i32) {addr_space = 0 : i32, dso_local} : i32
 @internal = internal global i32 42
-; CHECK: llvm.mlir.global available_externally @available_externally(42 : i32) : i32
+; CHECK: llvm.mlir.global available_externally @available_externally(42 : i32) {addr_space = 0 : i32}  : i32
 @available_externally = available_externally global i32 42
-; CHECK: llvm.mlir.global linkonce @linkonce(42 : i32) : i32
+; CHECK: llvm.mlir.global linkonce @linkonce(42 : i32) {addr_space = 0 : i32} : i32
 @linkonce = linkonce global i32 42
-; CHECK: llvm.mlir.global weak @weak(42 : i32) : i32
+; CHECK: llvm.mlir.global weak @weak(42 : i32) {addr_space = 0 : i32} : i32
 @weak = weak global i32 42
-; CHECK: llvm.mlir.global common @common(0 : i32) : i32
+; CHECK: llvm.mlir.global common @common(0 : i32) {addr_space = 0 : i32} : i32
 @common = common global i32 zeroinitializer
-; CHECK: llvm.mlir.global appending @appending(dense<[0, 1]> : tensor<2xi32>) : !llvm.array<2 x i32>
+; CHECK: llvm.mlir.global appending @appending(dense<[0, 1]> : tensor<2xi32>) {addr_space = 0 : i32} : !llvm.array<2 x i32>
 @appending = appending global [2 x i32] [i32 0, i32 1]
-; CHECK: llvm.mlir.global extern_weak @extern_weak() : i32
+; CHECK: llvm.mlir.global extern_weak @extern_weak() {addr_space = 0 : i32} : i32
 @extern_weak = extern_weak global i32
-; CHECK: llvm.mlir.global linkonce_odr @linkonce_odr(42 : i32) : i32
+; CHECK: llvm.mlir.global linkonce_odr @linkonce_odr(42 : i32) {addr_space = 0 : i32} : i32
 @linkonce_odr = linkonce_odr global i32 42
-; CHECK: llvm.mlir.global weak_odr @weak_odr(42 : i32) : i32
+; CHECK: llvm.mlir.global weak_odr @weak_odr(42 : i32) {addr_space = 0 : i32} : i32
 @weak_odr = weak_odr global i32 42
-; CHECK: llvm.mlir.global external @external() : i32
+; CHECK: llvm.mlir.global external @external() {addr_space = 0 : i32} : i32
 @external = external global i32
 
 ;
 ; UnnamedAddr attribute.
 ;
 
 
-; CHECK: llvm.mlir.global private constant @no_unnamed_addr(42 : i64) {dso_local} : i64
+; CHECK: llvm.mlir.global private constant @no_unnamed_addr(42 : i64) {addr_space = 0 : i32, dso_local} : i64
 @no_unnamed_addr = private constant i64 42
-; CHECK: llvm.mlir.global private local_unnamed_addr constant @local_unnamed_addr(42 : i64) {dso_local} : i64
+; CHECK: llvm.mlir.global private local_unnamed_addr constant @local_unnamed_addr(42 : i64) {addr_space = 0 : i32, dso_local} : i64
 @local_unnamed_addr = private local_unnamed_addr constant i64 42
-; CHECK: llvm.mlir.global private unnamed_addr constant @unnamed_addr(42 : i64) {dso_local} : i64
+; CHECK: llvm.mlir.global private unnamed_addr constant @unnamed_addr(42 : i64) {addr_space = 0 : i32, dso_local} : i64
 @unnamed_addr = private unnamed_addr constant i64 42
 
 ;
 ; Section attribute
 ;
 
-; CHECK: llvm.mlir.global internal constant @sectionvar("teststring") {dso_local, section = ".mysection"}
+; CHECK: llvm.mlir.global internal constant @sectionvar("teststring") {addr_space = 0 : i32, dso_local, section = ".mysection"}
 @sectionvar = internal constant [10 x i8] c"teststring", section ".mysection"
 
 ;
 ; Sequential constants.
 ;
 
-; CHECK: llvm.mlir.global internal constant @vector_constant(dense<[1, 2]> : vector<2xi32>) {dso_local} : vector<2xi32>
+; CHECK: llvm.mlir.global internal constant @vector_constant(dense<[1, 2]> : vector<2xi32>) {addr_space = 0 : i32, dso_local} : vector<2xi32>
 @vector_constant = internal constant <2 x i32> <i32 1, i32 2>
-; CHECK: llvm.mlir.global internal constant @array_constant(dense<[1.000000e+00, 2.000000e+00]> : tensor<2xf32>) {dso_local} : !llvm.array<2 x f32>
+; CHECK: llvm.mlir.global internal constant @array_constant(dense<[1.000000e+00, 2.000000e+00]> : tensor<2xf32>) {addr_space = 0 : i32, dso_local} : !llvm.array<2 x f32>
 @array_constant = internal constant [2 x float] [float 1., float 2.]
-; CHECK: llvm.mlir.global internal constant @nested_array_constant(dense<[{{\[}}1, 2], [3, 4]]> : tensor<2x2xi32>) {dso_local} : !llvm.array<2 x array<2 x i32>>
+; CHECK: llvm.mlir.global internal constant @nested_array_constant(dense<[{{\[}}1, 2], [3, 4]]> : tensor<2x2xi32>) {addr_space = 0 : i32, dso_local} : !llvm.array<2 x array<2 x i32>>
 @nested_array_constant = internal constant [2 x [2 x i32]] [[2 x i32] [i32 1, i32 2], [2 x i32] [i32 3, i32 4]]
-; CHECK: llvm.mlir.global internal constant @nested_array_constant3(dense<[{{\[}}[1, 2], [3, 4]]]> : tensor<1x2x2xi32>) {dso_local} : !llvm.array<1 x array<2 x array<2 x i32>>>
+; CHECK: llvm.mlir.global internal constant @nested_array_constant3(dense<[{{\[}}[1, 2], [3, 4]]]> : tensor<1x2x2xi32>) {addr_space = 0 : i32, dso_local} : !llvm.array<1 x array<2 x array<2 x i32>>>
 @nested_array_constant3 = internal constant [1 x [2 x [2 x i32]]] [[2 x [2 x i32]] [[2 x i32] [i32 1, i32 2], [2 x i32] [i32 3, i32 4]]]
-; CHECK: llvm.mlir.global internal constant @nested_array_vector(dense<[{{\[}}[1, 2], [3, 4]]]> : vector<1x2x2xi32>) {dso_local} : !llvm.array<1 x array<2 x vector<2xi32>>>
+; CHECK: llvm.mlir.global internal constant @nested_array_vector(dense<[{{\[}}[1, 2], [3, 4]]]> : vector<1x2x2xi32>) {addr_space = 0 : i32, dso_local} : !llvm.array<1 x array<2 x vector<2xi32>>>
 @nested_array_vector = internal constant [1 x [2 x <2 x i32>]] [[2 x <2 x i32>] [<2 x i32> <i32 1, i32 2>, <2 x i32> <i32 3, i32 4>]]
 
 ;
 ; Linkage on functions.
 ;
 
 ; CHECK: llvm.func internal @func_internal
 define internal void @func_internal() {

mlir/test/Target/LLVMIR/Import/zeroinitializer.ll

 ; RUN: mlir-translate --import-llvm %s | FileCheck %s
 
 %Domain = type { %Domain**, %Domain* }
 
-; CHECK: llvm.mlir.global external @D() :
-; CHECK-SAME: !llvm.struct<"Domain", (ptr<ptr<struct<"Domain">>>, ptr<struct<"Domain">>)> 
+; CHECK: llvm.mlir.global external @D()
+; CHECK-SAME: !llvm.struct<"Domain", (ptr<ptr<struct<"Domain">>>, ptr<struct<"Domain">>)>
 ; CHECK-DAG: %[[E0:.+]] = llvm.mlir.null : !llvm.ptr<struct<"Domain", (ptr<ptr<struct<"Domain">>>, ptr<struct<"Domain">>)>>
 ; CHECK-DAG: %[[E1:.+]] = llvm.mlir.null : !llvm.ptr<ptr<struct<"Domain", (ptr<ptr<struct<"Domain">>>, ptr<struct<"Domain">>)>>>
 ; CHECK: %[[ROOT:.+]] = llvm.mlir.undef : !llvm.struct<"Domain", (ptr<ptr<struct<"Domain">>>, ptr<struct<"Domain">>)>
 ; CHECK: %[[CHAIN:.+]] = llvm.insertvalue %[[E1]], %[[ROOT]][0 : i32]
 ; CHECK: %[[RES:.+]] = llvm.insertvalue %[[E0]], %[[CHAIN]][1 : i32]
 ; CHECK: llvm.return %[[RES]]
 @D = global %Domain zeroinitializer
 

mlir/test/mlir-linalg-ods-gen/test-linalg-ods-yaml-gen.yaml

 #       ODS:  let summary = [{ Title. }];
 #  ODS-NEXT:  let description = [{
 #  ODS-NEXT:    Detailed description.
 #  ODS-NEXT:  }];
 
 #       ODS:  let arguments =
 #  ODS-NEXT:    Variadic<AnyType>:$inputs,
 #  ODS-NEXT:    Variadic<AnyShaped>:$outputs,
-#  ODS-NEXT:    DefaultValuedAttr<TypeFnAttr, "TypeFn::cast_signed">:$cast
+#  ODS-NEXT:    DefaultValuedOptionalAttr<TypeFnAttr, "TypeFn::cast_signed">:$cast
 
 #       ODS:  let builders =
 #       ODS:  (ins "TypeRange":$resultTensorTypes, "ValueRange":$inputs,
 #  ODS-NEXT:       "ValueRange":$outputs,
 #  ODS-NEXT:       CArg<"ArrayRef<NamedAttribute>", "{}">:$attributes),
 
 #       ODS:  (ins "TypeRange":$resultTensorTypes, "ValueRange":$inputs,
 #  ODS-NEXT:       "ValueRange":$outputs, "Attribute":$cast,
- !ScalarAssign
     value: !ScalarExpression
       scalar_arg: I
 
 # ODS-LABEL:  def Test2Op : LinalgStructuredBase_Op<"test2"
 
 #       ODS:  let arguments =
 #  ODS-NEXT:    Variadic<AnyType>:$inputs,
 #  ODS-NEXT:    Variadic<AnyShaped>:$outputs,
-#  ODS-NEXT:    DefaultValuedAttr<RankedI64ElementsAttr<[2]>
+#  ODS-NEXT:    DefaultValuedOptionalAttr<RankedI64ElementsAttr<[2]>
 #  ODS-SAME:    "{ static_cast<int64_t>(1), static_cast<int64_t>(2) }">:$strides
 
 #       ODS:  "Attribute":$strides
 #       ODS:  $_state.addAttribute("strides", strides);
 
 #       ODS:  bool hasDynamicIndexingMaps();
 #  ODS-NEXT:  LogicalResult verifyIndexingMapRequiredAttributes();
 
value: !ScalarExpression
         - !ScalarExpression
           scalar_arg: O
 
 # ODS-LABEL:  def Test4Op : LinalgStructuredBase_Op<"test4"
 
 #       ODS:  let arguments =
 #  ODS-NEXT:    Variadic<AnyType>:$inputs,
 #  ODS-NEXT:    Variadic<AnyShaped>:$outputs,
-#  ODS-NEXT:    DefaultValuedAttr<UnaryFnAttr, "UnaryFn::exp">:$unary_fun,
-#  ODS-NEXT:    DefaultValuedAttr<BinaryFnAttr, "BinaryFn::add">:$binary_fun
+#  ODS-NEXT:    DefaultValuedOptionalAttr<UnaryFnAttr, "UnaryFn::exp">:$unary_fun,
+#  ODS-NEXT:    DefaultValuedOptionalAttr<BinaryFnAttr, "BinaryFn::add">:$binary_fun
 
 #       ODS:    "Attribute":$unary_fun, "Attribute":$binary_fun,
 
 #       ODS:    $_state.addAttribute("unary_fun", unary_fun)
 #  ODS-NEXT:    $_state.addAttribute("binary_fun", binary_fun)
 
 # IMPL-LABEL:  void Test4Op::regionBuilder(ImplicitLocOpBuilder &b,
 #  IMPL-NEXT:    Block &block, ArrayRef<NamedAttribute> attrs)

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

 // DEF:      some-return-type AOp::aAttr() {
 // DEF-NEXT:   auto attr = aAttrAttr()
 // DEF-NEXT:   return attr.some-convert-from-storage();
 
 // DEF:      some-attr-kind AOp::bAttrAttr()
 // DEF-NEXT:   ::mlir::impl::getAttrFromSortedRange((*this)->getAttrs().begin() + 1, (*this)->getAttrs().end() - 0, bAttrAttrName()).dyn_cast_or_null<some-attr-kind>()
 // DEF:      some-return-type AOp::bAttr() {
 // DEF-NEXT:   auto attr = bAttrAttr();
-// DEF-NEXT:   if (!attr)
-// DEF-NEXT:       return some-const-builder-call(::mlir::Builder((*this)->getContext()), 4.2).some-convert-from-storage();
 // DEF-NEXT:   return attr.some-convert-from-storage();
 
 // DEF:      some-attr-kind AOp::cAttrAttr()
 // DEF-NEXT:   ::mlir::impl::getAttrFromSortedRange((*this)->getAttrs().begin() + 1, (*this)->getAttrs().end() - 0, cAttrAttrName()).dyn_cast_or_null<some-attr-kind>()
 // DEF:      ::llvm::Optional<some-return-type> AOp::cAttr() {
 // DEF-NEXT:   auto attr = cAttrAttr()
 // DEF-NEXT:   return attr ? ::llvm::Optional<some-return-type>(attr.some-convert-from-storage()) : (::llvm::None);
 

 // DEF: ::mlir::ArrayAttr BOp::some_attr_array()
 // DEF: ::mlir::Type BOp::type_attr()
 
 // Test building constant values for array attribute kinds
 // ---
 
 def COp : NS_Op<"c_op", []> {
   let arguments = (ins
-    DefaultValuedAttr<I32ArrayAttr, "{1, 2}">:$i32_array_attr,
-    DefaultValuedAttr<I64ArrayAttr, "{3, 4}">:$i64_array_attr,
-    DefaultValuedAttr<F32ArrayAttr, "{5.f, 6.f}">:$f32_array_attr,
-    DefaultValuedAttr<F64ArrayAttr, "{7., 8.}">:$f64_array_attr,
-    DefaultValuedAttr<StrArrayAttr, "{\"a\", \"b\"}">:$str_array_attr
+    DefaultValuedOptionalAttr<I32ArrayAttr, "{1, 2}">:$i32_array_attr,
+    DefaultValuedOptionalAttr<I64ArrayAttr, "{3, 4}">:$i64_array_attr,
+    DefaultValuedOptionalAttr<F32ArrayAttr, "{5.f, 6.f}">:$f32_array_attr,
+    DefaultValuedOptionalAttr<F64ArrayAttr, "{7., 8.}">:$f64_array_attr,
+    DefaultValuedOptionalAttr<StrArrayAttr, "{\"a\", \"b\"}">:$str_array_attr
   );
 }
 
 // DEF-LABEL: COp definitions
 // DEF: ::mlir::Builder((*this)->getContext()).getI32ArrayAttr({1, 2})
 // DEF: ::mlir::Builder((*this)->getContext()).getI64ArrayAttr({3, 4})
 // DEF: ::mlir::Builder((*this)->getContext()).getF32ArrayAttr({5.f, 6.f})
 // DEF: ::mlir::Builder((*this)->getContext()).getF64ArrayAttr({7., 8.})

mlir/test/mlir-tblgen/op-python-bindings.td

 // CHECK: class DefaultValuedAttrsOp(_ods_ir.OpView):
 // CHECK-LABEL: OPERATION_NAME = "test.default_valued_attrs"
 def DefaultValuedAttrsOp : TestOp<"default_valued_attrs"> {
   // CHECK: def __init__(self, *, arr=None, unsupported=None, loc=None, ip=None):
   // CHECK:   operands = []
   // CHECK:   results = []
   // CHECK:   attributes = {}
   // CHECK:   regions = None
-  // CHECK:   attributes["arr"] = arr if arr is not None else _ods_ir.ArrayAttr.get([])
-  // CHECK:   unsupported is not None, "attribute unsupported must be specified"
+  // CHECK:   if arr is not None: attributes["arr"] = arr
+  // CHECK:   if unsupported is not None: attributes["unsupported"] = unsupported
   // CHECK:   _ods_successors = None
   // CHECK:   super().__init__(self.build_generic(
   // CHECK:     attributes=attributes, results=results, operands=operands,
   // CHECK:     successors=_ods_successors, regions=regions, loc=loc, ip=ip))
 
   let arguments = (ins DefaultValuedAttr<I64ArrayAttr, "{}">:$arr,
                        DefaultValuedAttr<I64ArrayAttr, "dummy_func()">:$unsupported);
   let results = (outs);

mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp

if (llvm::any_of(opConfig.structuredOp->args, [](LinalgOperandDef &arg) {
     for (LinalgOperandDef &arg : opConfig.structuredOp->args) {
       static const char paramFmt[] = "\"Attribute\":${0}";
       static const char stmtFmt[] = "$_state.addAttribute(\"{0}\", {0});";
       // Add the type conversion attributes to the op definition and builders.
       if (isFunctionAttribute(arg.kind)) {
         assert(arg.defaultFn);
         std::string enumName = convertOperandKindToEnumName(arg.kind);
         static const char typeFmt[] = "{0}::{1}";
-        static const char defFmt[] = "DefaultValuedAttr<{0}, \"{1}\">:${2}";
+        static const char defFmt[] = "DefaultValuedOptionalAttr<{0}, \"{1}\">:${2}";
         attrDefs.push_back(llvm::formatv(
             defFmt, llvm::formatv("{0}Attr", enumName),
             llvm::formatv(typeFmt, enumName, arg.defaultFn), arg.name));
         attrParams.push_back(llvm::formatv(paramFmt, arg.name));
         attrStmts.push_back(llvm::formatv(stmtFmt, arg.name));
       }
       // Add the index attributes to the op definition and builders.
       if (arg.kind == LinalgOperandDefKind::IndexAttr) {
         assert(arg.indexAttrMap.has_value());
         assert(arg.defaultIndices.has_value());
         size_t size = arg.indexAttrMap->affineMap().getNumResults();
         assert(arg.defaultIndices.value().size() == size);
         static const char typeFmt[] = "RankedI64ElementsAttr<[{0}]>";
-        static const char defFmt[] = "DefaultValuedAttr<{0}, \"{ {1} }\">:${2}";
+        static const char defFmt[] = "DefaultValuedOptionalAttr<{0}, \"{ {1} }\">:${2}";
         std::string defaultVals;
         llvm::raw_string_ostream ss(defaultVals);
         llvm::interleave(
             arg.defaultIndices.value(), ss,
             [&](int64_t val) { ss << "static_cast<int64_t>(" << val << ")"; },
             ", ");
         attrDefs.push_back(llvm::formatv(defFmt, llvm::formatv(typeFmt, size),
                                          ss.str(), arg.name));

mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp

 static void emitAttrGetterWithReturnType(FmtContext &fctx,
                                          OpClassOrAdaptor &opClass,
                                          const Operator &op, StringRef name,
                                          Attribute attr) {
   auto *method = opClass.addMethod(attr.getReturnType(), name);
   ERROR_IF_PRUNED(method, name, op);
   auto &body = method->body();
   body << "  auto attr = " << name << "Attr();\n";
-  if (attr.hasDefaultValue()) {
+  if (attr.hasDefaultValue() && attr.isOptional()) {
     // Returns the default value if not set.
     // TODO: this is inefficient, we are recreating the attribute for every
     // call. This should be set instead.
     if (!attr.isConstBuildable()) {
       PrintFatalError("DefaultValuedAttr of type " + attr.getAttrDefName() +
                       " must have a constBuilder");
     }
     std::string defaultValue = std::string(
void OpEmitter::genPopulateDefaultAttributes() {
   body << "::mlir::Builder " << odsBuilder
        << "(attrNames.front().getContext());\n";
   StringMap<int> attrIndex;
   for (const auto &it : llvm::enumerate(emitHelper.getAttrMetadata())) {
     attrIndex[it.value().first] = it.index();
   }
   for (const NamedAttribute &namedAttr : op.getAttributes()) {
     auto &attr = namedAttr.attr;
-    if (!attr.hasDefaultValue())
+    if (!attr.hasDefaultValue() || attr.isOptional())
       continue;
     auto index = attrIndex[namedAttr.name];
     body << "if (!attributes.get(attrNames[" << index << "])) {\n";
     FmtContext fctx;
     fctx.withBuilder(odsBuilder);
     std::string defaultValue = std::string(
         tgfmt(attr.getConstBuilderTemplate(), &fctx, attr.getDefaultValue()));
     body.indent() << formatv(" attributes.append(attrNames[{0}], {1});\n",
auto emitAttrWithStorageType = [&](StringRef name, StringRef emitName,
     auto &body = method->body().indent();
     body << "assert(odsAttrs && \"no attributes when constructing adapter\");\n"
          << formatv("auto attr = {0}.{1}<{2}>();\n", emitHelper.getAttr(name),
                     attr.hasDefaultValue() || attr.isOptional()
                         ? "dyn_cast_or_null"
                         : "cast",
                     attr.getStorageType());
 
-    if (attr.hasDefaultValue()) {
+    if (attr.hasDefaultValue() && attr.isOptional()) {
       // Use the default value if attribute is not set.
       // TODO: this is inefficient, we are recreating the attribute for every
       // call. This should be set instead.
       std::string defaultValue = std::string(
           tgfmt(attr.getConstBuilderTemplate(), &fctx, attr.getDefaultValue()));
       body << "  if (!attr)\n    attr = " << defaultValue << ";\n";
     }
     body << "return attr;\n";

mlir/tools/mlir-tblgen/OpPythonBindingGen.cpp

 //===- OpPythonBindingGen.cpp - Generator of Python API for MLIR Ops ------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // OpPythonBindingGen uses ODS specification of MLIR ops to generate Python
 // binding classes wrapping a generic operation API.
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/Support/LogicalResult.h"
 #include "mlir/TableGen/GenInfo.h"
 #include "mlir/TableGen/Operator.h"
 #include "llvm/ADT/StringSet.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/FormatVariadic.h"
 #include "llvm/TableGen/Error.h"
 #include "llvm/TableGen/Record.h"
 

 constexpr const char *initAttributeTemplate = R"Py(attributes["{0}"] = {1})Py";
 
 /// Template for setting an optional attribute in the operation builder.
 ///   {0} is the attribute name;
 ///   {1} is the builder argument name.
 constexpr const char *initOptionalAttributeTemplate =
     R"Py(if {1} is not None: attributes["{0}"] = {1})Py";
 
-/// Template for setting an attribute with a default value in the operation
-/// builder.
-///   {0} is the attribute name;
-///   {1} is the builder argument name;
-///   {2} is the default value.
-constexpr const char *initDefaultValuedAttributeTemplate =
-    R"Py(attributes["{0}"] = {1} if {1} is not None else {2})Py";
-
-/// Template for asserting that an attribute value was provided when calling a
-/// builder.
-///   {0} is the attribute name;
-///   {1} is the builder argument name.
-constexpr const char *assertAttributeValueSpecified =
-    R"Py(assert {1} is not None, "attribute {0} must be specified")Py";
-
 constexpr const char *initUnitAttributeTemplate =
     R"Py(if bool({1}): attributes["{0}"] = _ods_ir.UnitAttr.get(
       _ods_get_default_loc_context(loc)))Py";
 
 /// Template to initialize the successors list in the builder if there are any
 /// successors.
 ///   {0} is the value to initialize the successors list to.
 constexpr const char *initSuccessorsTemplate = R"Py(_ods_successors = {0})Py";
for (int i = 0, e = op.getNumSuccessors(); i < e; ++i) {
     if (name.empty())
       name = llvm::formatv("_gen_successor_{0}", i);
     name = sanitizeName(name);
     builderArgs.push_back(name);
     successorArgNames.push_back(name);
   }
 }
 
-/// Generates Python code for the default value of the given attribute.
-static FailureOr<std::string> getAttributeDefaultValue(Attribute attr) {
-  assert(attr.hasDefaultValue() && "expected attribute with default value");
-  StringRef storageType = attr.getStorageType().trim();
-  StringRef defaultValCpp = attr.getDefaultValue().trim();
-
-  // A list of commonly used attribute types and default values for which
-  // we can generate Python code. Extend as needed.
-  if (storageType.equals("::mlir::ArrayAttr") && defaultValCpp.equals("{}"))
-    return std::string("_ods_ir.ArrayAttr.get([])");
-
-  // No match: Cannot generate Python code.
-  return failure();
-}
-
 /// Populates `builderLines` with additional lines that are required in the
 /// builder to set up operation attributes. `argNames` is expected to contain
 /// the names of builder arguments that correspond to op arguments, i.e. to the
 /// operands and attributes in the same order as they appear in the `arguments`
 /// field.
 static void
 populateBuilderLinesAttr(const Operator &op,
                          llvm::ArrayRef<std::string> argNames,
                          llvm::SmallVectorImpl<std::string> &builderLines) {
   for (int i = 0, e = op.getNumArgs(); i < e; ++i) {
     Argument arg = op.getArg(i);
     auto *attribute = arg.dyn_cast<NamedAttribute *>();
     if (!attribute)
       continue;
 
     // Unit attributes are handled specially.
     if (attribute->attr.getStorageType().trim().equals("::mlir::UnitAttr")) {
       builderLines.push_back(llvm::formatv(initUnitAttributeTemplate,
                                            attribute->name, argNames[i]));
       continue;
     }
 
-    // Attributes with default value are handled specially.
-    if (attribute->attr.hasDefaultValue()) {
-      // In case we cannot generate Python code for the default value, the
-      // attribute must be specified by the user.
-      FailureOr<std::string> defaultValPy =
-          getAttributeDefaultValue(attribute->attr);
-      if (succeeded(defaultValPy)) {
-        builderLines.push_back(llvm::formatv(initDefaultValuedAttributeTemplate,
-                                             attribute->name, argNames[i],
-                                             *defaultValPy));
-      } else {
-        builderLines.push_back(llvm::formatv(assertAttributeValueSpecified,
-                                             attribute->name, argNames[i]));
-        builderLines.push_back(
-            llvm::formatv(initAttributeTemplate, attribute->name, argNames[i]));
-      }
-      continue;
-    }
-
-    builderLines.push_back(llvm::formatv(attribute->attr.isOptional()
-                                             ? initOptionalAttributeTemplate
-                                             : initAttributeTemplate,
-                                         attribute->name, argNames[i]));
+    builderLines.push_back(llvm::formatv(
+        (attribute->attr.isOptional() || attribute->attr.hasDefaultValue())
   }
 }
 
 /// Populates `builderLines` with additional lines that are required in the
 /// builder to set up successors. successorArgNames is expected to correspond
 /// to the Python argument name for each successor on the op.
 static void populateBuilderLinesSuccessors(
     const Operator &op, llvm::ArrayRef<std::string> successorArgNames,

mlir/unittests/IR/OperationSupportTest.cpp

 TEST(OperandStorageTest, PopulateDefaultAttrs) {
   MLIRContext context;
   context.getOrLoadDialect<test::TestDialect>();
   Builder builder(&context);
 
   OpBuilder b(&context);
   auto req1 = b.getI32IntegerAttr(10);
   auto req2 = b.getI32IntegerAttr(60);
+  // Verify default attributes populated post op creation.
   Operation *op = b.create<test::OpAttrMatch1>(b.getUnknownLoc(), req1, nullptr,
                                                nullptr, req2);
-  EXPECT_EQ(op->getAttr("default_valued_attr"), nullptr);
-  op->populateDefaultAttrs();
   auto opt = op->getAttr("default_valued_attr");
   EXPECT_NE(opt, nullptr) << *op;
 
   op->destroy();
 }
 } // namespace