[Mlir-commits] [mlir] 7d273fd - [mlir] Populate default attributes on op creation
Jacques Pienaar
llvmlistbot at llvm.org
Mon Aug 22 16:49:58 PDT 2022
Author: Jacques Pienaar
Date: 2022-08-22T16:49:46-07:00
New Revision: 7d273fde110d7735512c3b71a83eb88e89d189cc
URL: https://github.com/llvm/llvm-project/commit/7d273fde110d7735512c3b71a83eb88e89d189cc
DIFF: https://github.com/llvm/llvm-project/commit/7d273fde110d7735512c3b71a83eb88e89d189cc.diff
LOG: [mlir] Populate default attributes on op creation
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 some of the OpenMP ones to optional attribute with default as the
currently set default values are not legal. May need to dig more there.
Switched LinAlg generated ones to optional attribute with default as its
quite widely used and unclear where it falls on two different
interpretations.
Differential Revision: https://reviews.llvm.org/D130916
Added:
Modified:
flang/test/Fir/convert-to-llvm.fir
flang/test/Fir/external-mangling.fir
flang/test/Fir/global-initialization.fir
mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
mlir/lib/IR/Operation.cpp
mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir
mlir/test/Dialect/GPU/outlining.mlir
mlir/test/Dialect/LLVMIR/global.mlir
mlir/test/Dialect/Linalg/vectorize-convolution.mlir
mlir/test/Target/LLVMIR/Import/basic.ll
mlir/test/Target/LLVMIR/Import/zeroinitializer.ll
mlir/test/mlir-linalg-ods-gen/test-linalg-ods-yaml-gen.yaml
mlir/test/mlir-tblgen/op-attribute.td
mlir/test/mlir-tblgen/op-python-bindings.td
mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp
mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp
mlir/tools/mlir-tblgen/OpPythonBindingGen.cpp
mlir/unittests/IR/OperationSupportTest.cpp
Removed:
################################################################################
diff --git a/flang/test/Fir/convert-to-llvm.fir b/flang/test/Fir/convert-to-llvm.fir
index 2bd27ee121543..252d7c669e71e 100644
--- a/flang/test/Fir/convert-to-llvm.fir
+++ b/flang/test/Fir/convert-to-llvm.fir
@@ -14,7 +14,7 @@ fir.global @g_i0 : i32 {
fir.has_value %1 : i32
}
-// CHECK: llvm.mlir.global external @g_i0() : i32 {
+// CHECK: llvm.mlir.global external @g_i0() {addr_space = 0 : i32} : i32 {
// CHECK: %[[C0:.*]] = llvm.mlir.constant(0 : i32) : i32
// CHECK: llvm.return %[[C0]] : i32
// CHECK: }
@@ -26,7 +26,7 @@ fir.global @g_ci5 constant : i32 {
fir.has_value %c : i32
}
-// CHECK: llvm.mlir.global external constant @g_ci5() : i32 {
+// CHECK: llvm.mlir.global external constant @g_ci5() {addr_space = 0 : i32} : i32 {
// CHECK: %[[C5:.*]] = llvm.mlir.constant(5 : i32) : i32
// CHECK: llvm.return %[[C5]] : i32
// CHECK: }
@@ -34,22 +34,22 @@ fir.global @g_ci5 constant : i32 {
// -----
fir.global internal @i_i515 (515:i32) : i32
-// CHECK: llvm.mlir.global internal @i_i515(515 : i32) : i32
+// CHECK: llvm.mlir.global internal @i_i515(515 : i32) {addr_space = 0 : i32} : i32
// -----
fir.global common @C_i511 (0:i32) : i32
-// CHECK: llvm.mlir.global common @C_i511(0 : i32) : i32
+// CHECK: llvm.mlir.global common @C_i511(0 : i32) {addr_space = 0 : i32} : i32
// -----
fir.global weak @w_i86 (86:i32) : i32
-// CHECK: llvm.mlir.global weak @w_i86(86 : i32) : i32
+// CHECK: llvm.mlir.global weak @w_i86(86 : i32) {addr_space = 0 : i32} : i32
// -----
fir.global linkonce @w_i86 (86:i32) : i32
-// CHECK: llvm.mlir.global linkonce @w_i86(86 : i32) : i32
+// CHECK: llvm.mlir.global linkonce @w_i86(86 : i32) {addr_space = 0 : i32} : i32
// -----
@@ -67,7 +67,7 @@ fir.global @symbol : i64 {
// CHECK: %{{.*}} = llvm.mlir.addressof @[[SYMBOL:.*]] : !llvm.ptr<i64>
-// CHECK: llvm.mlir.global external @[[SYMBOL]]() : i64 {
+// CHECK: llvm.mlir.global external @[[SYMBOL]]() {addr_space = 0 : i32} : i64 {
// CHECK: %{{.*}} = llvm.mlir.constant(1 : i64) : i64
// CHECK: llvm.return %{{.*}} : i64
// CHECK: }
@@ -84,7 +84,7 @@ fir.global internal @_QEmultiarray : !fir.array<32x32xi32> {
fir.has_value %2 : !fir.array<32x32xi32>
}
-// CHECK: llvm.mlir.global internal @_QEmultiarray() : !llvm.array<32 x array<32 x i32>> {
+// CHECK: llvm.mlir.global internal @_QEmultiarray() {addr_space = 0 : i32} : !llvm.array<32 x array<32 x i32>> {
// CHECK: %[[CST:.*]] = llvm.mlir.constant(dense<1> : vector<32x32xi32>) : !llvm.array<32 x array<32 x i32>>
// CHECK: llvm.return %[[CST]] : !llvm.array<32 x array<32 x i32>>
// CHECK: }
@@ -101,7 +101,7 @@ fir.global internal @_QEmultiarray : !fir.array<32xi32> {
fir.has_value %2 : !fir.array<32xi32>
}
-// CHECK: llvm.mlir.global internal @_QEmultiarray() : !llvm.array<32 x i32> {
+// CHECK: llvm.mlir.global internal @_QEmultiarray() {addr_space = 0 : i32} : !llvm.array<32 x i32> {
// CHECK: %[[CST:.*]] = llvm.mlir.constant(1 : i32) : i32
// CHECK: %{{.*}} = llvm.mlir.undef : !llvm.array<32 x i32>
// CHECK: %{{.*}} = llvm.insertvalue %[[CST]], %{{.*}}[5] : !llvm.array<32 x i32>
@@ -1635,7 +1635,7 @@ func.func @embox1(%arg0: !fir.ref<!fir.type<_QMtest_dinitTtseq{i:i32}>>) {
return
}
-// CHECK: llvm.mlir.global linkonce constant @_QMtest_dinitE.dt.tseq() : i8
+// CHECK: llvm.mlir.global linkonce constant @_QMtest_dinitE.dt.tseq() {addr_space = 0 : i32} : i8
// CHECK-LABEL: llvm.func @embox1
// CHECK: %[[TYPE_CODE:.*]] = llvm.mlir.constant(42 : i32) : i32
// CHECK: %[[TYPE_CODE_I8:.*]] = llvm.trunc %[[TYPE_CODE]] : i32 to i8
diff --git a/flang/test/Fir/external-mangling.fir b/flang/test/Fir/external-mangling.fir
index 9e26e96eac9d5..794132847c6de 100644
--- a/flang/test/Fir/external-mangling.fir
+++ b/flang/test/Fir/external-mangling.fir
@@ -2,7 +2,7 @@
// RUN: tco --external-name-interop %s | FileCheck %s
// RUN: tco --external-name-interop %s | tco --fir-to-llvm-ir | FileCheck %s --check-prefix=LLVMIR
-func.func @_QPfoo() {
+func.func @_QPfoo() {
%c0 = arith.constant 0 : index
%0 = fir.address_of(@_QBa) : !fir.ref<!fir.array<4xi8>>
%1 = fir.convert %0 : (!fir.ref<!fir.array<4xi8>>) -> !fir.ref<!fir.array<?xi8>>
@@ -35,7 +35,7 @@ func.func private @_QPbar2(!fir.ref<f32>)
// LLVMIR: llvm.call @bar_(%{{.*}}) : (!llvm.ptr<i32>) -> ()
// LLVMIR: llvm.call @bar2_(%{{.*}}) : (!llvm.ptr<f32>) -> ()
-// LLVMIR: llvm.mlir.global common @a_(dense<0> : vector<4xi8>) : !llvm.array<4 x i8>
-// LLVMIR: llvm.mlir.global common @__BLNK__(dense<0> : vector<4xi8>) : !llvm.array<4 x i8>
+// LLVMIR: llvm.mlir.global common @a_(dense<0> : vector<4xi8>) {{.*}} : !llvm.array<4 x i8>
+// LLVMIR: llvm.mlir.global common @__BLNK__(dense<0> : vector<4xi8>) {{.*}} : !llvm.array<4 x i8>
// LLVMIR: llvm.func @bar_(!llvm.ptr<i32>) attributes {sym_visibility = "private"}
// LLVMIR: llvm.func @bar2_(!llvm.ptr<f32>) attributes {sym_visibility = "private"}
diff --git a/flang/test/Fir/global-initialization.fir b/flang/test/Fir/global-initialization.fir
index 78340fa25b0e4..28361a9438adc 100644
--- a/flang/test/Fir/global-initialization.fir
+++ b/flang/test/Fir/global-initialization.fir
@@ -8,7 +8,7 @@ fir.global internal @_QEmask : !fir.array<32xi32> {
fir.has_value %2 : !fir.array<32xi32>
}
-// CHECK: llvm.mlir.global internal @_QEmask() : !llvm.array<32 x i32> {
+// CHECK: llvm.mlir.global internal @_QEmask() {addr_space = 0 : i32} : !llvm.array<32 x i32> {
// CHECK: [[VAL0:%.*]] = llvm.mlir.constant(1 : i32) : i32
// CHECK: [[VAL1:%.*]] = llvm.mlir.undef : !llvm.array<32 x i32>
// CHECK: [[VAL2:%.*]] = llvm.mlir.constant(dense<1> : vector<32xi32>) : !llvm.array<32 x i32>
@@ -22,7 +22,7 @@ fir.global internal @_QEmultiarray : !fir.array<32x32xi32> {
fir.has_value %2 : !fir.array<32x32xi32>
}
-// CHECK: llvm.mlir.global internal @_QEmultiarray() : !llvm.array<32 x array<32 x i32>> {
+// CHECK: llvm.mlir.global internal @_QEmultiarray() {addr_space = 0 : i32} : !llvm.array<32 x array<32 x i32>> {
// CHECK: [[VAL0:%.*]] = llvm.mlir.constant(1 : i32) : i32
// CHECK: [[VAL1:%.*]] = llvm.mlir.undef : !llvm.array<32 x array<32 x i32>>
// CHECK: [[VAL2:%.*]] = llvm.mlir.constant(dense<1> : vector<32x32xi32>) : !llvm.array<32 x array<32 x i32>>
@@ -37,7 +37,7 @@ fir.global internal @_QEmasklogical : !fir.array<32768x!fir.logical<4>> {
fir.has_value %2 : !fir.array<32768x!fir.logical<4>>
}
-// CHECK: llvm.mlir.global internal @_QEmasklogical() : !llvm.array<32768 x i32> {
+// CHECK: llvm.mlir.global internal @_QEmasklogical() {addr_space = 0 : i32} : !llvm.array<32768 x i32> {
// CHECK: [[VAL0:%.*]] = llvm.mlir.constant(true) : i1
// CHECK: [[VAL1:%.*]] = llvm.mlir.undef : !llvm.array<32768 x i32>
// CHECK: [[VAL2:%.*]] = llvm.zext [[VAL0]] : i1 to i32
@@ -57,7 +57,7 @@ fir.global internal @_QElookforme : !fir.type<_QTt{i:!fir.array<500xi32>,j:!fir.
fir.has_value %5 : !fir.type<_QTt{i:!fir.array<500xi32>,j:!fir.array<500xi32>}>
}
-// CHECK: llvm.mlir.global internal @_QElookforme() : !llvm.struct<"_QTt", (array<500 x i32>, array<500 x i32>)> {
+// CHECK: llvm.mlir.global internal @_QElookforme() {addr_space = 0 : i32} : !llvm.struct<"_QTt", (array<500 x i32>, array<500 x i32>)> {
// CHECK: [[CST0:%.*]] = llvm.mlir.constant(2 : i32) : i32
// CHECK: [[CST1:%.*]] = llvm.mlir.constant(52 : i32) : i32
// CHECK: [[STRUCT:%.*]] = llvm.mlir.undef : !llvm.struct<"_QTt", (array<500 x i32>, array<500 x i32>)>
diff --git a/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td b/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
index 54673332b65a7..1ddb5ea31d397 100644
--- a/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
+++ b/mlir/include/mlir/Dialect/OpenMP/OpenMPOps.td
@@ -1046,7 +1046,7 @@ def AtomicReadOp : OpenMP_Op<"atomic.read", [AllTypesMatch<["x", "v"]>]> {
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
@@ -1093,7 +1093,7 @@ def AtomicWriteOp : OpenMP_Op<"atomic.write"> {
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
@@ -1157,7 +1157,7 @@ def AtomicUpdateOp : OpenMP_Op<"atomic.update",
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 = [{
@@ -1219,7 +1219,7 @@ def AtomicCaptureOp : OpenMP_Op<"atomic.capture",
}];
- 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 = [{
diff --git a/mlir/lib/IR/Operation.cpp b/mlir/lib/IR/Operation.cpp
index e30d5ea1ad235..886b5bb58c925 100644
--- a/mlir/lib/IR/Operation.cpp
+++ b/mlir/lib/IR/Operation.cpp
@@ -77,6 +77,10 @@ Operation *Operation::create(Location location, OperationName name,
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,
diff --git a/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir b/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir
index 0a50960f7903b..f6e496c713f03 100644
--- a/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir
+++ b/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir
@@ -615,16 +615,16 @@ func.func @transpose(%arg0: memref<?x?x?xf32, offset: ?, strides: [?, ?, 1]>) {
// -----
-// 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.
@@ -672,7 +672,7 @@ func.func @get_gv2_memref() {
}
// 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
@@ -691,7 +691,7 @@ func.func @get_gv3_memref() {
}
// 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}
// -----
diff --git a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir b/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
index 2ac727f4c28d3..984b5184bbab9 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
+++ b/mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
@@ -2,7 +2,7 @@
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
diff --git a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
index ea68dc9d57189..095a3f11a54d3 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
+++ b/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
@@ -30,12 +30,12 @@ spv.func @access_chain_array(%arg0 : i32) "None" {
//===----------------------------------------------------------------------===//
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>
@@ -46,7 +46,7 @@ spv.module Logical GLSL450 {
}
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>
@@ -57,7 +57,7 @@ spv.module Logical GLSL450 {
}
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>
@@ -68,7 +68,7 @@ spv.module @name Logical GLSL450 {
}
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" {
@@ -78,7 +78,7 @@ spv.module Logical GLSL450 {
}
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" {
diff --git a/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir
index 54420fca6080f..e438e03ebb375 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir
+++ b/mlir/test/Conversion/SPIRVToLLVM/misc-ops-to-llvm.mlir
@@ -89,7 +89,7 @@ spv.func @vector_shuffle_
diff erent_size(%vector1: vector<3xf32>, %vector2: vecto
//===----------------------------------------------------------------------===//
// 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] : !llvm.struct<(i32)>
@@ -108,7 +108,7 @@ spv.module Logical OpenCL {
}
// 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] : !llvm.struct<(i32, array<3 x i32>)>
@@ -120,7 +120,7 @@ spv.module Logical OpenCL {
// CHECK-NEXT: %[[RET:.*]] = llvm.insertvalue %[[C2]], %[[T2]][1, 2] : !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: }
diff --git a/mlir/test/Dialect/GPU/outlining.mlir b/mlir/test/Dialect/GPU/outlining.mlir
index 89080c5871fbc..5191dcf8fffb8 100644
--- a/mlir/test/Dialect/GPU/outlining.mlir
+++ b/mlir/test/Dialect/GPU/outlining.mlir
@@ -286,7 +286,7 @@ func.func @recursive_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()
diff --git a/mlir/test/Dialect/LLVMIR/global.mlir b/mlir/test/Dialect/LLVMIR/global.mlir
index 9454662235691..37c55ecfe7a0d 100644
--- a/mlir/test/Dialect/LLVMIR/global.mlir
+++ b/mlir/test/Dialect/LLVMIR/global.mlir
@@ -6,16 +6,16 @@ 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")
@@ -27,7 +27,7 @@ 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: }
@@ -57,9 +57,9 @@ llvm.mlir.global extern_weak @extern_weak() : i64
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
@@ -81,13 +81,13 @@ func.func @references() {
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>
// -----
diff --git a/mlir/test/Dialect/Linalg/vectorize-convolution.mlir b/mlir/test/Dialect/Linalg/vectorize-convolution.mlir
index 7e1f39cbda3e9..e3c4d7d1c1de0 100644
--- a/mlir/test/Dialect/Linalg/vectorize-convolution.mlir
+++ b/mlir/test/Dialect/Linalg/vectorize-convolution.mlir
@@ -38,14 +38,14 @@ func.func @conv1d_nwc_4x2x8_memref(%input: memref<4x6x3xf32>, %filter: memref<1x
/// 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>
@@ -104,25 +104,25 @@ func.func @conv1d_nwc_4x2x8_memref(%input: memref<4x6x3xf32>, %filter: memref<2x
/// 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>
@@ -172,13 +172,13 @@ func.func @conv1d_nwc_4x2x8_memref(%input: memref<4x6x3xf32>, %filter: memref<2x
/// 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>
diff --git a/mlir/test/Target/LLVMIR/Import/basic.ll b/mlir/test/Target/LLVMIR/Import/basic.ll
index 9f5f8bb8ffbc2..113ba62bb4cad 100644
--- a/mlir/test/Target/LLVMIR/Import/basic.ll
+++ b/mlir/test/Target/LLVMIR/Import/basic.ll
@@ -3,24 +3,24 @@
%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>
@@ -32,14 +32,14 @@
; 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
;
@@ -53,27 +53,27 @@
; 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
;
@@ -81,33 +81,33 @@
;
-; 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>]]
;
diff --git a/mlir/test/Target/LLVMIR/Import/zeroinitializer.ll b/mlir/test/Target/LLVMIR/Import/zeroinitializer.ll
index b30a54aafb28e..3f582138b03b5 100644
--- a/mlir/test/Target/LLVMIR/Import/zeroinitializer.ll
+++ b/mlir/test/Target/LLVMIR/Import/zeroinitializer.ll
@@ -2,7 +2,7 @@
%Domain = type { %Domain**, %Domain* }
-; CHECK: llvm.mlir.global external @D() :
+; 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">>)>>>
diff --git a/mlir/test/mlir-linalg-ods-gen/test-linalg-ods-yaml-gen.yaml b/mlir/test/mlir-linalg-ods-gen/test-linalg-ods-yaml-gen.yaml
index 23cccc7e484db..81f40df090066 100644
--- a/mlir/test/mlir-linalg-ods-gen/test-linalg-ods-yaml-gen.yaml
+++ b/mlir/test/mlir-linalg-ods-gen/test-linalg-ods-yaml-gen.yaml
@@ -70,7 +70,7 @@ structured_op: !LinalgStructuredOpConfig
# 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,
@@ -157,7 +157,7 @@ structured_op: !LinalgStructuredOpConfig
# 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
@@ -305,8 +305,8 @@ structured_op: !LinalgStructuredOpConfig
# 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,
diff --git a/mlir/test/mlir-tblgen/op-attribute.td b/mlir/test/mlir-tblgen/op-attribute.td
index 6c55fc45d38c9..3c85d71460dc5 100644
--- a/mlir/test/mlir-tblgen/op-attribute.td
+++ b/mlir/test/mlir-tblgen/op-attribute.td
@@ -109,8 +109,6 @@ def AOp : NS_Op<"a_op", []> {
// 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()
@@ -343,11 +341,11 @@ def BOp : NS_Op<"b_op", []> {
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
);
}
diff --git a/mlir/test/mlir-tblgen/op-python-bindings.td b/mlir/test/mlir-tblgen/op-python-bindings.td
index 2b73132a7431d..2dda3db53bb2e 100644
--- a/mlir/test/mlir-tblgen/op-python-bindings.td
+++ b/mlir/test/mlir-tblgen/op-python-bindings.td
@@ -188,8 +188,8 @@ def DefaultValuedAttrsOp : TestOp<"default_valued_attrs"> {
// 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,
diff --git a/mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp b/mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp
index 9f88176713b1f..56bd8968f660f 100644
--- a/mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp
+++ b/mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp
@@ -722,7 +722,8 @@ static LogicalResult generateNamedGenericOpOds(LinalgOpConfig &opConfig,
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));
@@ -736,7 +737,8 @@ static LogicalResult generateNamedGenericOpOds(LinalgOpConfig &opConfig,
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(
diff --git a/mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp b/mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp
index c341591f807cc..c3a06df700c0a 100644
--- a/mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp
+++ b/mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp
@@ -952,7 +952,7 @@ static void emitAttrGetterWithReturnType(FmtContext &fctx,
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.
@@ -1611,7 +1611,7 @@ void OpEmitter::genPopulateDefaultAttributes() {
}
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";
@@ -2912,7 +2912,7 @@ OpOperandAdaptorEmitter::OpOperandAdaptorEmitter(
: "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.
diff --git a/mlir/tools/mlir-tblgen/OpPythonBindingGen.cpp b/mlir/tools/mlir-tblgen/OpPythonBindingGen.cpp
index e40d0ff8faf21..f83908c368383 100644
--- a/mlir/tools/mlir-tblgen/OpPythonBindingGen.cpp
+++ b/mlir/tools/mlir-tblgen/OpPythonBindingGen.cpp
@@ -11,7 +11,6 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Support/LogicalResult.h"
#include "mlir/TableGen/GenInfo.h"
#include "mlir/TableGen/Operator.h"
#include "llvm/ADT/StringSet.h"
@@ -543,21 +542,6 @@ constexpr const char *initAttributeTemplate = R"Py(attributes["{0}"] = {1})Py";
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";
@@ -663,21 +647,6 @@ static void populateBuilderArgsSuccessors(
}
}
-/// 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
@@ -700,29 +669,11 @@ populateBuilderLinesAttr(const Operator &op,
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())
+ ? initOptionalAttributeTemplate
+ : initAttributeTemplate,
+ attribute->name, argNames[i]));
}
}
diff --git a/mlir/unittests/IR/OperationSupportTest.cpp b/mlir/unittests/IR/OperationSupportTest.cpp
index 598ef9ed156a5..ff9ab385893b7 100644
--- a/mlir/unittests/IR/OperationSupportTest.cpp
+++ b/mlir/unittests/IR/OperationSupportTest.cpp
@@ -281,10 +281,9 @@ TEST(OperandStorageTest, PopulateDefaultAttrs) {
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;
More information about the Mlir-commits
mailing list