Integrate LLVM at llvm/llvm-project@1cbf8e89b54d
Updates LLVM usage to match
[1cbf8e89b54d](https://github.com/llvm/llvm-project/commit/1cbf8e89b54d)
PiperOrigin-RevId: 341867974
diff --git a/SUBMODULE_VERSIONS b/SUBMODULE_VERSIONS
index a1dc4c9..38e0fdd 100644
--- a/SUBMODULE_VERSIONS
+++ b/SUBMODULE_VERSIONS
@@ -6,7 +6,7 @@
4fb0ff7069bd88ee85902f4d0bb62794e5f6d021 third_party/flatcc
f2fb48c3b3d79a75a88a99fba6576b25d42ec528 third_party/googletest
e3f662922d8340a7fd46a4f75cad0e72163bf219 third_party/llvm-bazel
-f147f59cd377a6be68e5ca5c343eb11df8e7ee6f third_party/llvm-project
+1cbf8e89b54de939420d53d7a528bec6fbaf0a55 third_party/llvm-project
17b12a4481daa150e2d1ea3ada086b551b856707 third_party/marl
9e970e18057e80e8af07f96d73714e395bac98d1 third_party/mlir-emitc
d8c7ee00a687ac369e62e2032514a93a9b413502 third_party/pybind11
diff --git a/iree/compiler/Conversion/HLOToLinalg/test/pipeline_test.mlir b/iree/compiler/Conversion/HLOToLinalg/test/pipeline_test.mlir
index 402166f..15c4ec5 100644
--- a/iree/compiler/Conversion/HLOToLinalg/test/pipeline_test.mlir
+++ b/iree/compiler/Conversion/HLOToLinalg/test/pipeline_test.mlir
@@ -87,7 +87,7 @@
// CHECK-DAG: %[[ARG1:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<512xf32>
// CHECK-DAG: %[[RET0:.+]] = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<4x512xf32>
// CHECK: linalg.generic
-// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]], %[[ARG1]], %[[ARG1]] :
+// CHECK-SAME: ins(%[[ARG0]], %[[ARG1]] :
// CHECK-SAME: outs(%[[RET0]] :
// -----
diff --git a/iree/compiler/Conversion/HLOToLinalg/test/pw_linalg_fusion.mlir b/iree/compiler/Conversion/HLOToLinalg/test/pw_linalg_fusion.mlir
deleted file mode 100644
index b65cf46..0000000
--- a/iree/compiler/Conversion/HLOToLinalg/test/pw_linalg_fusion.mlir
+++ /dev/null
@@ -1,76 +0,0 @@
-// RUN: iree-opt -split-input-file -iree-codegen-hlo-to-linalg-on-tensors -linalg-fusion-for-tensor-ops %s | IreeFileCheck %s
-
-// CHECK-LABEL: @pw_fusion_two
-func @pw_fusion_two(%arg0: tensor<4x8xi32>, %arg1: tensor<4x8xi32>, %arg2 : tensor<4x8xi32>) -> tensor<4x8xi32> {
- // CHECK: linalg.generic
- // CHECK: ^{{[a-zA-Z0-9$._-]+}}
- // CHECK-SAME: %[[ARG0:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG2:[a-zA-Z0-9$._-]+]]: i32
- // CHECK: %[[TEMP:[a-zA-Z0-9$._-]+]] = muli %[[ARG0]], %[[ARG1]]
- // CHECK: addi %[[TEMP]], %[[ARG2]]
- // CHECK-NOT: linalg.generic
- %4 = "mhlo.multiply"(%arg0, %arg1) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- %5 = "mhlo.add"(%4, %arg2) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- return %5 : tensor<4x8xi32>
-}
-
-// -----
-
-// CHECK-LABEL: @pw_fusion_three
-func @pw_fusion_three(%arg0: tensor<4x8xi32>, %arg1: tensor<4x8xi32>, %arg2 : tensor<4x8xi32>, %arg3: tensor<4x8xi32>) -> tensor<4x8xi32> {
- // CHECK: linalg.generic
- // CHECK: ^{{[a-zA-Z0-9$._-]+}}
- // CHECK-SAME: %[[ARG0:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG2:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG3:[a-zA-Z0-9$._-]+]]: i32
- // CHECK: %[[TEMP1:[a-zA-Z0-9$._-]+]] = muli %[[ARG0]], %[[ARG1]]
- // CHECK: %[[TEMP2:[a-zA-Z0-9$._-]+]] = addi %[[TEMP1]], %[[ARG2]]
- // CHECK: subi %[[TEMP2]], %[[ARG3]]
- // CHECK-NOT: linalg.generic
- %4 = "mhlo.multiply"(%arg0, %arg1) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- %5 = "mhlo.add"(%4, %arg2) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- %6 = "mhlo.subtract"(%5, %arg3) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- return %6: tensor<4x8xi32>
-}
-
-// -----
-
-// CHECK-LABEL: @pw_fusion_dag
-func @pw_fusion_dag(%arg0: tensor<4x8xi32>, %arg1: tensor<4x8xi32>, %arg2 : tensor<4x8xi32>, %arg3: tensor<4x8xi32>) -> tensor<4x8xi32> {
- // CHECK: linalg.generic
- // CHECK: ^{{[a-zA-Z0-9$._-]+}}
- // CHECK-SAME: %[[ARG0:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG2:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG3:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-DAG: %[[TEMP1:[a-zA-Z0-9$._-]+]] = muli %[[ARG0]], %[[ARG1]]
- // CHECK-DAG: %[[TEMP2:[a-zA-Z0-9$._-]+]] = addi %[[ARG2]], %[[ARG3]]
- // CHECK: subi %[[TEMP1]], %[[TEMP2]]
- // CHECK-NOT: linalg.generic
- %4 = "mhlo.multiply"(%arg0, %arg1) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- %5 = "mhlo.add"(%arg2, %arg3) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- %6 = "mhlo.subtract"(%4, %5) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- return %6: tensor<4x8xi32>
-}
-
-// -----
-
-// CHECK-LABEL: @pw_fusion_dag2
-func @pw_fusion_dag2(%arg0: tensor<4x8xi32>, %arg1: tensor<4x8xi32>, %arg2 : tensor<4x8xi32>) -> tensor<4x8xi32> {
- // CHECK: linalg.generic
- // CHECK: ^{{[a-zA-Z0-9$._-]+}}
- // CHECK-SAME: %[[ARG0:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG1:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG2:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-SAME: %[[ARG3:[a-zA-Z0-9$._-]+]]: i32
- // CHECK-DAG: %[[TEMP1:[a-zA-Z0-9$._-]+]] = muli %[[ARG0]], %[[ARG1]]
- // CHECK-DAG: %[[TEMP2:[a-zA-Z0-9$._-]+]] = addi %[[ARG2]], %[[ARG3]]
- // CHECK: subi %[[TEMP1]], %[[TEMP2]]
- // CHECK-NOT: linalg.generic
- %3 = "mhlo.multiply"(%arg0, %arg1) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- %4 = "mhlo.add"(%arg0, %arg2) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- %5 = "mhlo.subtract"(%3, %4) : (tensor<4x8xi32>, tensor<4x8xi32>) -> tensor<4x8xi32>
- return %5: tensor<4x8xi32>
-}
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir
index c4419d8..a63c723 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir
+++ b/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir
@@ -126,7 +126,7 @@
// CHECK: load %{{.+}}[%[[IV0]], %[[IV1]], %[[IV2]], %[[IV3]]]
// CHECK: store %{{.+}}[%[[IV0]], %[[IV1]], %[[IV2]], %[[IV3]]]
-// CHECK: func @[[NUM_WORKGROUPS_FN]]
+// CHECK: func private @[[NUM_WORKGROUPS_FN]]
// CHECK-DAG: %[[C1:.+]] = constant 1 : index
// CHECK-DAG: %[[C12:.+]] = constant 12 : index
// CHECK: return %[[C12]], %[[C1]], %[[C1]]
@@ -179,7 +179,7 @@
// CHECK-NEXT: store
// CHECK-NEXT: return
-// CHECK: func @[[NUM_WORKGROUPS_FN]]
+// CHECK: func private @[[NUM_WORKGROUPS_FN]]
// CHECK-DAG: %[[C1:.+]] = constant 1 : index
// CHECK: return %[[C1]], %[[C1]], %[[C1]]
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/linalg_tile_and_fuse.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/linalg_tile_and_fuse.mlir
index 1f27ed9..a115363 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/linalg_tile_and_fuse.mlir
+++ b/iree/compiler/Conversion/LinalgToSPIRV/test/linalg_tile_and_fuse.mlir
@@ -84,7 +84,7 @@
// CHECK: linalg.conv
// CHECK-SAME: %[[ARG0]], %[[VIEW1]], %[[VIEW2]]
// CHECK-SAME: "workgroup"
-// CHECK: func @[[NUM_WORKGROUPS_FN]]
+// CHECK: func private @[[NUM_WORKGROUPS_FN]]
// CHECK-DAG: %[[C0:.+]] = constant 0
// CHECK-DAG: %[[C1:.+]] = constant 1
// CHECK-DAG: %[[C2:.+]] = constant 2
@@ -158,7 +158,7 @@
// CHECK-SAME: "workgroup"
// CHECK-SAME: ins(%[[VIEW0]], %[[VIEW1]]
// CHECK-SAME: outs(%[[VIEW2]]
-// CHECK: func @[[NUM_WORKGROUPS_FN]]
+// CHECK: func private @[[NUM_WORKGROUPS_FN]]
// CHECK-DAG: %[[C8:.+]] = constant 8 : index
// CHECK-DAG: %[[C7:.+]] = constant 7 : index
// CHECK-DAG: %[[C0:.+]] = constant 0 : index
@@ -222,7 +222,7 @@
// CHECK: linalg.pooling_max
// CHECK-SAME: %[[VIEW0]], %[[ARG1]], %[[VIEW2]]
// CHECK-SAME: "workgroup"
-// CHECK: func @[[NUM_WORKGROUPS_FN]]
+// CHECK: func private @[[NUM_WORKGROUPS_FN]]
// CHECK-DAG: %[[C0:.+]] = constant 0
// CHECK-DAG: %[[C1:.+]] = constant 1
// CHECK-DAG: %[[C32:.+]] = constant 32
@@ -288,7 +288,7 @@
// CHECK: linalg.pooling_max
// CHECK-SAME: %[[VIEW0]], %[[ARG1]], %[[VIEW2]]
// CHECK-SAME: "workgroup"
-// CHECK: func @[[NUM_WORKGROUPS_FN]]
+// CHECK: func private @[[NUM_WORKGROUPS_FN]]
// CHECK-DAG: %[[C1:.+]] = constant 1
// CHECK-DAG: %[[C32:.+]] = constant 32
// CHECK-DAG: %[[C31:.+]] = constant 31
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/split_dispatch_function.mlir b/iree/compiler/Conversion/LinalgToSPIRV/test/split_dispatch_function.mlir
index c543614..7f7ad16 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/split_dispatch_function.mlir
+++ b/iree/compiler/Conversion/LinalgToSPIRV/test/split_dispatch_function.mlir
@@ -166,7 +166,7 @@
// CHECK: linalg.conv(%[[IN2]], %[[TS1]], %[[TS2]])
// CHECK: return
-// CHECK: func @[[NUM_WORKGROUPS_FN2]]
+// CHECK: func private @[[NUM_WORKGROUPS_FN2]]
// CHECK: func @kernel_dispatch_1()
// CHECK-SAME: {hal.num_workgroups_fn = @[[NUM_WORKGROUPS_FN1:.+]]}
@@ -176,7 +176,7 @@
// CHECK: scf.yield
// CHECK: return
-// CHECK: func @[[NUM_WORKGROUPS_FN1]]
+// CHECK: func private @[[NUM_WORKGROUPS_FN1]]
// CHECK: func @kernel_dispatch_0()
// CHECK-SAME: {hal.num_workgroups_fn = @[[NUM_WORKGROUPS_FN0:.+]]}
@@ -188,7 +188,7 @@
// CHECK: linalg.fill(%[[TS]], %[[ZERO]])
// CHECK: return
-// CHECK: func @[[NUM_WORKGROUPS_FN0]]
+// CHECK: func private @[[NUM_WORKGROUPS_FN0]]
func @kernel() attributes {hal.num_workgroups_fn = @kernel__num_workgroups__} {
%cst = constant 0.000000e+00 : f32
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/identify_constant_pools.mlir b/iree/compiler/Dialect/HAL/Transforms/test/identify_constant_pools.mlir
index 3c58678..56a3474 100644
--- a/iree/compiler/Dialect/HAL/Transforms/test/identify_constant_pools.mlir
+++ b/iree/compiler/Dialect/HAL/Transforms/test/identify_constant_pools.mlir
@@ -27,7 +27,7 @@
// CHECK: flow.variable @variable_0 mutable init(@variable_0_initializer)
flow.variable @variable_0 mutable dense<3.0> : tensor<128xf32>
-// CHECK-NEXT: func @variable_0_initializer() -> tensor<128xf32>
+// CHECK-NEXT: func private @variable_0_initializer() -> tensor<128xf32>
// CHECK-NEXT: [[CONST:%.+]] = hal.constant_pool.load @_const_pool_init::@variable_0 : tensor<128xf32>
// CHECK-NEXT: return [[CONST]] : tensor<128xf32>
// CHECK-NEXT: }
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/materialize_constant_pool_buffers.mlir b/iree/compiler/Dialect/HAL/Transforms/test/materialize_constant_pool_buffers.mlir
index 54b5ae2..3232083 100644
--- a/iree/compiler/Dialect/HAL/Transforms/test/materialize_constant_pool_buffers.mlir
+++ b/iree/compiler/Dialect/HAL/Transforms/test/materialize_constant_pool_buffers.mlir
@@ -10,7 +10,7 @@
}
// CHECK: hal.variable @dense_variable_init_storage_buffer init(@dense_variable_init_storage_buffer_initializer) : !hal.buffer
-// CHECK-NEXT: func @dense_variable_init_storage_buffer_initializer() -> !hal.buffer
+// CHECK-NEXT: func private @dense_variable_init_storage_buffer_initializer() -> !hal.buffer
// CHECK: [[STORAGE:%.+]] = hal.constant_storage.lookup @dense_variable_init::@_storage : !iree.byte_buffer
// CHECK: = hal.allocator.map {{.+}} [[STORAGE]][%c0, %c768] : !iree.byte_buffer -> !hal.buffer
@@ -25,7 +25,7 @@
}
// CHECK: hal.variable @splat_variable_init_splats init(@splat_variable_init_splats_initializer) : !hal.buffer
-// CHECK-NEXT: func @splat_variable_init_splats_initializer() -> !hal.buffer
+// CHECK-NEXT: func private @splat_variable_init_splats_initializer() -> !hal.buffer
// CHECK: [[BUFFER:%.+]] = hal.allocator.allocate {{.+}} %c64 : !hal.buffer
// CHECK: hal.buffer.fill [[BUFFER]], %c0, %c4, %c1065353216_i32
// CHECK: hal.buffer.fill [[BUFFER]], %c32, %c32_0, %c1234567890_i32
@@ -47,15 +47,15 @@
}
// CHECK: hal.variable @pool_storage0_buffer init(@pool_storage0_buffer_initializer) : !hal.buffer
-// CHECK-NEXT: func @pool_storage0_buffer_initializer() -> !hal.buffer
+// CHECK-NEXT: func private @pool_storage0_buffer_initializer() -> !hal.buffer
// CHECK: [[STORAGE:%.+]] = hal.constant_storage.lookup @pool::@_storage0 : !iree.byte_buffer
// CHECK: = hal.allocator.map {{.+}} [[STORAGE]][%c0, %c16] : !iree.byte_buffer -> !hal.buffer
// CHECK: hal.variable @pool_storage1_buffer init(@pool_storage1_buffer_initializer) : !hal.buffer
-// CHECK-NEXT: func @pool_storage1_buffer_initializer() -> !hal.buffer
+// CHECK-NEXT: func private @pool_storage1_buffer_initializer() -> !hal.buffer
// CHECK: hal.variable @pool_splats init(@pool_splats_initializer) : !hal.buffer
-// CHECK-NEXT: func @pool_splats_initializer() -> !hal.buffer
+// CHECK-NEXT: func private @pool_splats_initializer() -> !hal.buffer
// CHECK: [[BUFFER:%.+]] = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %c64 : !hal.buffer
// CHECK: hal.buffer.fill [[BUFFER]], %c0, %c4, %c1065353216_i32
// CHECK: hal.buffer.fill [[BUFFER]], %c32, %c32_0, %c1234567890_i32
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir b/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
index 499d36f..e2640e7 100644
--- a/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
+++ b/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
@@ -24,7 +24,7 @@
// CHECK-NEXT: hal.interface.store.tensor %[[RET0]], @legacy_io::@ret0, offset = %[[ZERO]] : tensor<4xf32>
// CHECK-NEXT: return
// CHECK-NEXT: }
- // CHECK-NEXT: func @simpleMath_rgn_dispatch_0_impl
+ // CHECK-NEXT: func private @simpleMath_rgn_dispatch_0_impl
func @simpleMath_rgn_dispatch_0(%arg0: tensor<4xf32>) -> tensor<4xf32> {
%0 = mhlo.add %arg0, %arg0 : tensor<4xf32>
return %0 : tensor<4xf32>
@@ -64,7 +64,7 @@
// CHECK-NEXT: hal.interface.store.tensor %[[RET0_I8]], @legacy_io::@ret0, offset = %[[ZERO]] : tensor<4xi8>
// CHECK-NEXT: return
// CHECK-NEXT: }
- // CHECK-NEXT: func @bools_rgn_dispatch_0_impl(%arg0: tensor<4xi1>, %arg1: tensor<4xi1>) -> tensor<4xi1>
+ // CHECK-NEXT: func private @bools_rgn_dispatch_0_impl(%arg0: tensor<4xi1>, %arg1: tensor<4xi1>) -> tensor<4xi1>
func @bools_rgn_dispatch_0(%arg0: tensor<4xi1>, %arg1: tensor<4xi1>) -> tensor<4xi1> {
%0 = mhlo.and %arg0, %arg1 : tensor<4xi1>
%c = mhlo.constant dense<[false, false, true, false]> : tensor<4xi1>
@@ -96,7 +96,7 @@
// CHECK-NEXT: hal.interface.store.tensor %[[RET0]], @legacy_io::@ret0, offset = %[[ZERO]] : tensor<7x?x10xf32>
// CHECK-NEXT: return
// CHECK-NEXT: }
- // CHECK-NEXT: func @entry_impl
+ // CHECK-NEXT: func private @entry_impl
func @entry(%arg0: tensor<?x7x10xf32>, %arg1: index, %arg2: index) -> tensor<7x?x10xf32> {
%0 = shapex.make_ranked_shape %arg1 : (index) -> !shapex.ranked_shape<[?,7,10]>
%1 = shapex.make_ranked_shape %arg2 : (index) -> !shapex.ranked_shape<[7,?,10]>
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/materialize_resource_caches.mlir b/iree/compiler/Dialect/HAL/Transforms/test/materialize_resource_caches.mlir
index 590df15..aff893c 100644
--- a/iree/compiler/Dialect/HAL/Transforms/test/materialize_resource_caches.mlir
+++ b/iree/compiler/Dialect/HAL/Transforms/test/materialize_resource_caches.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -iree-hal-materialize-resource-caches %s -iree-hal-target-backends=vmla | IreeFileCheck %s
// CHECK: hal.variable @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !hal.descriptor_set_layout
-// CHECK-NEXT: func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
+// CHECK-NEXT: func private @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout {
// CHECK-NEXT: %dev = hal.ex.shared_device : !hal.device
// CHECK-NEXT: %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write">] : !hal.descriptor_set_layout
// CHECK-NEXT: return %descriptor_set_layout : !hal.descriptor_set_layout
@@ -23,7 +23,7 @@
// CHECK: hal.variable @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !hal.descriptor_set_layout
// CHECK: hal.variable @_executable_layout_0 init(@_executable_layout_0_initializer) : !hal.executable_layout
-// CHECK-NEXT: func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
+// CHECK-NEXT: func private @_executable_layout_0_initializer() -> !hal.executable_layout {
// CHECK-NEXT: %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
// CHECK-NEXT: %dev = hal.ex.shared_device : !hal.device
// CHECK-NEXT: %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
@@ -49,7 +49,7 @@
// CHECK: hal.variable @_descriptor_set_layout_1
// CHECK: hal.variable @_executable_layout_0 init(@_executable_layout_0_initializer) : !hal.executable_layout
-// CHECK-NEXT: func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
+// CHECK-NEXT: func private @_executable_layout_0_initializer() -> !hal.executable_layout {
// CHECK-NEXT: %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
// CHECK-NEXT: %1 = hal.variable.load @_descriptor_set_layout_1 : !hal.descriptor_set_layout
// CHECK-NEXT: %dev = hal.ex.shared_device : !hal.device
@@ -119,7 +119,7 @@
}
// CHECK: hal.variable @_executable_cache init(@_executable_cache_initializer) : !hal.executable_cache
-// CHECK-NEXT: func @_executable_cache_initializer
+// CHECK-NEXT: func private @_executable_cache_initializer
// CHECK: %[[CACHE:.+]] = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache
// CHECK-NEXT: hal.device.switch(%dev : !hal.device)
// CHECK-NEXT: #hal.device.match.id<"vmla">(%[[CACHE_CAPTURE:.+]] = %executable_cache_default : !hal.executable_cache) {
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/memoize_device_queries.mlir b/iree/compiler/Dialect/HAL/Transforms/test/memoize_device_queries.mlir
index a081a4c..b6b7339 100644
--- a/iree/compiler/Dialect/HAL/Transforms/test/memoize_device_queries.mlir
+++ b/iree/compiler/Dialect/HAL/Transforms/test/memoize_device_queries.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -iree-hal-memoize-device-queries %s | IreeFileCheck %s
// CHECK: hal.variable @_device_match_id_0 init(@_device_match_id_0_initializer) : i1
-// CHECK: func @_device_match_id_0_initializer() -> i1
+// CHECK: func private @_device_match_id_0_initializer() -> i1
// CHECK-NEXT: %[[DEVICE:.+]] = hal.ex.shared_device : !hal.device
// CHECK-NEXT: %[[IS_MATCH:.+]] = hal.device.match.id %[[DEVICE]], pattern = ["vulkan-v1.?-*"] : (!hal.device) -> i1
// CHECK-NEXT: return %[[IS_MATCH]] : i1
diff --git a/iree/compiler/Dialect/VM/Transforms/test/mark_public_symbols_exported.mlir b/iree/compiler/Dialect/VM/Transforms/test/mark_public_symbols_exported.mlir
index 8bd5256..3310761 100644
--- a/iree/compiler/Dialect/VM/Transforms/test/mark_public_symbols_exported.mlir
+++ b/iree/compiler/Dialect/VM/Transforms/test/mark_public_symbols_exported.mlir
@@ -1,8 +1,7 @@
// RUN: iree-opt -split-input-file -iree-vm-mark-public-symbols-exported %s | IreeFileCheck %s
-// CHECK-LABEL: @private_symbol
-// CHECK-SAME: {sym_visibility = "private"}
-func @private_symbol() attributes {sym_visibility = "private"}
+// CHECK-LABEL: private @private_symbol
+func private @private_symbol()
// CHECK-LABEL: @public_symbol
// CHECK-SAME: {iree.module.export}
diff --git a/iree/compiler/Dialect/VMLA/Conversion/HLOToVMLA/test/convert.mlir b/iree/compiler/Dialect/VMLA/Conversion/HLOToVMLA/test/convert.mlir
index 95759c6..8abb0f5 100644
--- a/iree/compiler/Dialect/VMLA/Conversion/HLOToVMLA/test/convert.mlir
+++ b/iree/compiler/Dialect/VMLA/Conversion/HLOToVMLA/test/convert.mlir
@@ -1,14 +1,14 @@
// RUN: iree-opt -split-input-file -iree-vmla-conversion %s | IreeFileCheck %s
-// CHECK-LABEL: func @basic
-func @basic(%arg0 : tensor<5xf32>) -> (tensor<5xi32>) attributes { sym_visibility = "private" } {
+// CHECK-LABEL: func private @basic
+func private @basic(%arg0 : tensor<5xf32>) -> (tensor<5xi32>) {
// CHECK: vmla.convert
%0 = "mhlo.convert"(%arg0) : (tensor<5xf32>) -> tensor<5xi32>
return %0 : tensor<5xi32>
}
-// CHECK-LABEL: func @noop
-func @noop(%arg0 : tensor<?xf32>) -> (tensor<5xf32>) attributes { sym_visibility = "private" } {
+// CHECK-LABEL: func private @noop
+func private @noop(%arg0 : tensor<?xf32>) -> (tensor<5xf32>) {
// CHECK: return %arg0
%0 = "mhlo.convert"(%arg0) : (tensor<?xf32>) -> tensor<5xf32>
return %0 : tensor<5xf32>
diff --git a/iree/compiler/Dialect/VMLA/Conversion/HLOToVMLA/test/sort.mlir b/iree/compiler/Dialect/VMLA/Conversion/HLOToVMLA/test/sort.mlir
index 0903793..826cfc8 100644
--- a/iree/compiler/Dialect/VMLA/Conversion/HLOToVMLA/test/sort.mlir
+++ b/iree/compiler/Dialect/VMLA/Conversion/HLOToVMLA/test/sort.mlir
@@ -1,6 +1,6 @@
// RUN: iree-opt -split-input-file -iree-vmla-pre-conversion-lowering -iree-vmla-conversion -canonicalize %s | IreeFileCheck %s
-func @sort1D(%arg0 : tensor<4xf32>) -> tensor<4xf32> attributes { sym_visibility = "private" } {
+func private @sort1D(%arg0 : tensor<4xf32>) -> tensor<4xf32> {
// CHECK-DAG: [[C16:%.+]] = constant 16 : index
// CHECK-DAG: [[RS:%.+]] = shapex.const_ranked_shape : !shapex.ranked_shape<[4]>
// CHECK-DAG: [[BL:%.+]] = vmla.buffer.alloc byte_length = [[C16]] : !vmla.buffer
@@ -18,8 +18,8 @@
}
-// CHECK-LABEL: func @sort2D
-func @sort2D(%arg0 : tensor<4x4xf32>) -> tensor<4x4xf32> attributes { sym_visibility = "private" } {
+// CHECK-LABEL: func private @sort2D
+func private @sort2D(%arg0 : tensor<4x4xf32>) -> tensor<4x4xf32> {
// CHECK-DAG: [[C64:%.+]] = constant 64 : index
// CHECK-DAG: [[RS:%.+]] = shapex.const_ranked_shape : !shapex.ranked_shape<[4,4]>
// CHECK-DAG: [[BL:%.+]] = vmla.buffer.alloc byte_length = [[C64]] : !vmla.buffer
diff --git a/iree/compiler/Dialect/VMLA/Transforms/test/pre_conversion_lowering.mlir b/iree/compiler/Dialect/VMLA/Transforms/test/pre_conversion_lowering.mlir
index 7d9e5ee..366dae2 100644
--- a/iree/compiler/Dialect/VMLA/Transforms/test/pre_conversion_lowering.mlir
+++ b/iree/compiler/Dialect/VMLA/Transforms/test/pre_conversion_lowering.mlir
@@ -16,8 +16,8 @@
// -----
-// CHECK-LABEL: func @f
-func @f(%arg0 : tensor<4xf32>) -> tensor<4xf32> attributes { sym_visibility = "private" } {
+// CHECK-LABEL: func private @f
+func private @f(%arg0 : tensor<4xf32>) -> tensor<4xf32> {
// CHECK-DAG: [[SORT:%.+]] = vmla.sort.pseudo %arg0
// CHECK-DAG: [[GATHER:%.+]] = "mhlo.torch_index_select"(%arg0, [[SORT]]) {batch_dims = 0 : i64, dim = 0 : i64}
%sort = "mhlo.sort"(%arg0) ( {
@@ -32,8 +32,8 @@
// -----
-// CHECK-LABEL: func @f
-func @f(%arg0 : tensor<4x4xf32>) -> tensor<4x4xf32> attributes { sym_visibility = "private" } {
+// CHECK-LABEL: func private @f
+func private @f(%arg0 : tensor<4x4xf32>) -> tensor<4x4xf32> {
// CHECK-DAG: [[SORT:%.+]] = vmla.sort.pseudo %arg0
// CHECK-DAG: [[GATHER:%.+]] = "mhlo.torch_index_select"(%arg0, [[SORT]]) {batch_dims = 1 : i64, dim = 1 : i64}
%sort = "mhlo.sort"(%arg0) ( {
@@ -66,8 +66,8 @@
// -----
-// CHECK-LABEL: func @f
-func @f(%arg0: tensor<8xcomplex<f32>>) -> tensor<8xcomplex<f32>> attributes { sym_visibility = "private" } {
+// CHECK-LABEL: func private @f
+func private @f(%arg0: tensor<8xcomplex<f32>>) -> tensor<8xcomplex<f32>> {
// CHECK-DAG: [[REAL:%.+]] = "mhlo.real"(%arg0)
// CHECK-DAG: [[IMAG:%.+]] = "mhlo.imag"(%arg0)
// CHECK-DAG: [[REAL_OUT:%.+]], [[IMAG_OUT:%.+]] = vmla.fft.pseudo [[REAL]], [[IMAG]]
