Synchronize submodules with LLVM at llvm/llvm-project@b9db70369b77 Updates LLVM dependencies to match [b9db70369b77](https://github.com/llvm/llvm-project/commit/b9db70369b77). - TensorFlow to [d7b3c9fad000](https://github.com/tensorflow/tensorflow/commit/d7b3c9fad000) - MLIR-HLO to [41d38dd3cce5](https://github.com/tensorflow/mlir-hlo/commit/${MLIR_HLO_SHA?}) `./scripts/git/update_to_llvm_syncpoint.py` Automated submodule bump from .github/workflows/update_llvm_dependent_submodules.yml PiperOrigin-RevId: 394310736
diff --git a/SUBMODULE_VERSIONS.txt b/SUBMODULE_VERSIONS.txt index e38c906..16a4f83 100644 --- a/SUBMODULE_VERSIONS.txt +++ b/SUBMODULE_VERSIONS.txt
@@ -4,7 +4,7 @@ aa533abfd4232b01f9e57041d70114d5a77e6de0 third_party/googletest 88b845dee001723c4a0db1fe5477de735b6d3bb0 third_party/liburing acd6f6f014c25e46363e718381e0b35205df2d83 third_party/libyaml -b9db70369b7799887b817e13109801795e4d70fc third_party/llvm-project +9b6c8132d3785269512803ff51cb421f8d8bcf0e third_party/llvm-project 41d38dd3cce59a1c02d3c0fde5b272d20952748c third_party/mlir-hlo 3f701faace7addc75d16dea8a6cd769fa5b3f260 third_party/musl 4c7697dbe973ed01ae6fbec37d186ebd05982e1f third_party/pybind11
diff --git a/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir b/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir index 30e8718..83d4865 100644 --- a/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir +++ b/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir
@@ -105,7 +105,7 @@ func @use_subspan_with_unrealized_conversion_cast(%offset : index, %i: index) -> f32 { %subspan = hal.interface.binding.subspan @io::@s0b0_ro_constant[%offset] : memref<6x7x8xf32> - %use = unrealized_conversion_cast %subspan : memref<6x7x8xf32> to memref<?xf32> + %use = builtin.unrealized_conversion_cast %subspan : memref<6x7x8xf32> to memref<?xf32> %val = memref.load %use[%i] : memref<?xf32> return %val: f32 }
diff --git a/iree/compiler/Codegen/Common/test/forop_canonicalization.mlir b/iree/compiler/Codegen/Common/test/forop_canonicalization.mlir index d106866..a619273 100644 --- a/iree/compiler/Codegen/Common/test/forop_canonicalization.mlir +++ b/iree/compiler/Codegen/Common/test/forop_canonicalization.mlir
@@ -35,19 +35,19 @@ %c0 = constant 0 : index %c1 = constant 1 : index %c10 = constant 10 : index - %0 = unrealized_conversion_cast %arg0 : vector<4xf32> to vector<1x4xf32> - %1 = unrealized_conversion_cast %arg1 : vector<4xf32> to vector<1x4xf32> + %0 = builtin.unrealized_conversion_cast %arg0 : vector<4xf32> to vector<1x4xf32> + %1 = builtin.unrealized_conversion_cast %arg1 : vector<4xf32> to vector<1x4xf32> %20:2 = scf.for %arg3 = %c0 to %c10 step %c1 iter_args(%arg4 = %0, %arg5 = %1) -> (vector<1x4xf32>, vector<1x4xf32>) { - %a = unrealized_conversion_cast %arg4 : vector<1x4xf32> to vector<4xf32> - %b = unrealized_conversion_cast %arg5 : vector<1x4xf32> to vector<4xf32> + %a = builtin.unrealized_conversion_cast %arg4 : vector<1x4xf32> to vector<4xf32> + %b = builtin.unrealized_conversion_cast %arg5 : vector<1x4xf32> to vector<4xf32> %c = addf %a, %b : vector<4xf32> %d = mulf %a, %b : vector<4xf32> - %cc = unrealized_conversion_cast %c : vector<4xf32> to vector<1x4xf32> - %dc = unrealized_conversion_cast %d : vector<4xf32> to vector<1x4xf32> + %cc = builtin.unrealized_conversion_cast %c : vector<4xf32> to vector<1x4xf32> + %dc = builtin.unrealized_conversion_cast %d : vector<4xf32> to vector<1x4xf32> scf.yield %cc, %dc : vector<1x4xf32>, vector<1x4xf32> } - %21 = unrealized_conversion_cast %20#0 : vector<1x4xf32> to vector<4xf32> - %22 = unrealized_conversion_cast %20#1 : vector<1x4xf32> to vector<4xf32> + %21 = builtin.unrealized_conversion_cast %20#0 : vector<1x4xf32> to vector<4xf32> + %22 = builtin.unrealized_conversion_cast %20#1 : vector<1x4xf32> to vector<4xf32> return %21, %22 : vector<4xf32>, vector<4xf32> }
diff --git a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir index c82a04d..98fba38 100644 --- a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir +++ b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
@@ -15,7 +15,7 @@ interface = @io, ordinal = 0 : index } - module { + builtin.module { func @matmul_tensors() { %c0 = constant 0 : index %c1 = constant 1 : index @@ -92,7 +92,7 @@ interface = @io, ordinal = 0 : index } - module { + builtin.module { func @add_no_config() { %c0 = constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<?x?xf32> @@ -136,7 +136,7 @@ interface = @io, ordinal = 0 : index, signature = (!flow.dispatch.tensor<readonly:?x?xf32>, !flow.dispatch.tensor<readonly:?xf32>, !flow.dispatch.tensor<writeonly:?x?xf32>) -> ()} - module { + builtin.module { func @add() { %c0 = constant 0 : index %c1 = constant 1 : index @@ -215,7 +215,7 @@ interface = @io, ordinal = 0 : index, signature = (!flow.dispatch.tensor<readonly:?x?x?x?xf32>, !flow.dispatch.tensor<readonly:?xf32>, !flow.dispatch.tensor<writeonly:?x?x?x?xf32>) -> ()} - module { + builtin.module { func @add4D() { %c0 = constant 0 : index %c1 = constant 1 : index @@ -315,7 +315,7 @@ interface = @io, ordinal = 0 : index } - module { + builtin.module { func @batch_matmul_tensors() { %c0 = constant 0 : index %c1 = constant 1 : index @@ -412,7 +412,7 @@ %13 = affine.min affine_map<(d0)[s0] -> (-d0 + 128, s0)>(%arg0)[%workgroup_size_y] %14 = affine.min affine_map<(d0)[s0] -> (-d0 + 512, s0)>(%arg1)[%workgroup_size_x] %15 = linalg.init_tensor [%13, %14] : tensor<?x?xf32> - %16 = linalg.fill(%cst, %15) : f32, tensor<?x?xf32> -> tensor<?x?xf32> + %16 = linalg.fill(%cst, %15) : f32, tensor<?x?xf32> -> tensor<?x?xf32> %17 = linalg.matmul {__internal_linalg_transform__ = "workgroup", lowering.config = {passPipeline = 1 : i32, tileSizes = [[32, 32, 32]]}} ins(%8, %10 : tensor<?x256xf32>, tensor<256x?xf32>) outs(%16 : tensor<?x?xf32>) -> tensor<?x?xf32> flow.dispatch.tensor.store %17, %2, offsets = [%arg0, %arg1], sizes = [%11, %12], strides = [1, 1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:128x512xf32> } @@ -438,7 +438,7 @@ // CHECK-DAG: %[[NWG_Y:.+]] = affine.apply #[[MAP0]]()[%[[ARG1]]] // CHECK: return %[[NWG_X]], %[[NWG_Y]], %[[C1]] // CHECK: builtin.module -// CHECK: builtin.func @preset_config +// CHECK: func @preset_config // CHECK-DAG: %[[WGID_X:.+]] = hal.interface.workgroup.id[0] // CHECK-DAG: %[[WGCOUNT_X:.+]] = hal.interface.workgroup.count[0] // CHECK-DAG: %[[WGID_Y:.+]] = hal.interface.workgroup.id[1]
diff --git a/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir b/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir index bcbd154..5bbb7ee 100644 --- a/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir +++ b/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
@@ -13,7 +13,7 @@ interface = @io, ordinal = 0 : index } - module { + builtin.module { func @matmul_128x128x128() { %c0 = constant 0 : index %c128 = constant 128 : index @@ -126,7 +126,7 @@ interface = @io, ordinal = 0 : index } - module { + builtin.module { func @matmul_i8_i8_i32_128x128x128() { %c0 = constant 0 : index %c128 = constant 128 : index
diff --git a/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir b/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir index 5a65678..4bbcb4b 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/distribute_wg_copy.mlir
@@ -17,7 +17,7 @@ memref.global "private" @__shared_memory___1 : memref<3x512xf32, 3> memref.global "private" @__shared_memory___0 : memref<256x4xf32, 3> memref.global "private" @__shared_memory__ : memref<64x16xf32, 3> - // CHECK-LABEL: builtin.func @shared_mem_cpy( + // CHECK-LABEL: @shared_mem_cpy( builtin.func @shared_mem_cpy( %m0 : memref<64x16xf32>, %m1 : memref<256x4xf32>, %m2 : memref<3x512xf32>) { %sm0 = memref.get_global @__shared_memory__ : memref<64x16xf32, 3>
diff --git a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir index 7034e05..0a02226 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
@@ -7,7 +7,7 @@ } hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @add_dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { func @add_dispatch_0() { %c0 = constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:16384xf32> @@ -52,7 +52,7 @@ hal.executable @dot_dispatch_1 attributes {sym_visibility = "private"} { hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @dot_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : index} - module { + builtin.module { func @dot_dispatch_1() { %c0 = constant 0 : index %c4 = constant 4 : index @@ -117,7 +117,7 @@ hal.executable.entry_point @predict_dispatch_153 attributes { interface = @io, ordinal = 0 : index} - module { + builtin.module { func @predict_dispatch_153() { %c0 = constant 0 : index %cst = constant 0x7FC00000 : f32
diff --git a/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir index e8122b7..d6025c4 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
@@ -10,7 +10,7 @@ } hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @add_dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { func @add_dispatch_0() { %c0 = constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:16xf32> @@ -53,7 +53,7 @@ } hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @dot_dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { func @dot_dispatch_0() { %cst = constant 0.000000e+00 : f32 %c0 = constant 0 : index @@ -141,7 +141,7 @@ } hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @dot_dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { func @dot_dispatch_0() { %cst = constant 0.000000e+00 : f32 %c0 = constant 0 : index @@ -209,7 +209,7 @@ hal.interface.binding @wo2, set=0, binding=2, type="StorageBuffer", access="Write|Discard" } hal.executable.entry_point @conv2d_dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { func @conv2d_dispatch_0() { %c0 = constant 0 : index %cst = constant 0.000000e+00 : f32 @@ -282,7 +282,7 @@ } hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @add_dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { func @add_dispatch_0() { %c0 = constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:16xf32> @@ -316,7 +316,7 @@ hal.executable @reduction_dispatch { hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @reduction attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { func @reduction() { %c0 = constant 0 : index %cst = constant 0.000000e+00 : f32
diff --git a/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir b/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir index a1c8079..6612b5f 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir
@@ -1,6 +1,6 @@ // RUN: iree-opt -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-llvmgpu-remove-single-iteration-loop))))' %s | IreeFileCheck %s -// CHECK-LABEL: builtin.func @dispatch_0() +// CHECK-LABEL: func @dispatch_0() hal.executable @dispatch_0 attributes {sym_visibility = "private"} { hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> { hal.executable.entry_point @dispatch_0 attributes {
diff --git a/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir index a4d3af2..8d4ec2a 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
@@ -10,7 +10,7 @@ } hal.executable.variant @rocm, target = #hal.executable.target<"rocm", "rocm-hsaco-fb"> { hal.executable.entry_point @add_dispatch_0 attributes {interface = @io, ordinal = 0 : index, signature = (!flow.dispatch.tensor<readonly:16xf32>, !flow.dispatch.tensor<readonly:16xf32>, !flow.dispatch.tensor<writeonly:16xf32>) -> ()} - module { + builtin.module { func @add_dispatch_0() { %c0 = constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:16xf32> @@ -53,7 +53,7 @@ } hal.executable.variant @rocm, target = #hal.executable.target<"rocm", "rocm-hsaco-fb"> { hal.executable.entry_point @dot_dispatch_0 attributes {interface = @io, ordinal = 0 : index, signature = (!flow.dispatch.tensor<readonly:1024x1024xf32>, !flow.dispatch.tensor<readonly:1024x1024xf32>, !flow.dispatch.tensor<writeonly:1024x1024xf32>) -> ()} - module { + builtin.module { func @dot_dispatch_0() { %cst = constant 0.000000e+00 : f32 %c0 = constant 0 : index
diff --git a/iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir b/iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir index cb01537..6d1cd27 100644 --- a/iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir +++ b/iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir
@@ -14,7 +14,7 @@ } return } -// CHECK-LABEL: builtin.func @add_dispatch_0 +// CHECK-LABEL: func @add_dispatch_0 // CHECK-COUNT-8: vector.transfer_read {{.*}} : memref<1x8x4xf32>, vector<1x1x4xf32> // CHECK-COUNT-8: vector.transfer_read {{.*}} : memref<1x4x8xf32>, vector<1x1x4xf32> // CHECK-COUNT-8: addf %{{.*}}, %{{.*}} : vector<1x1x4xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir b/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir index 18c5fdc..eb46b1c 100644 --- a/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir +++ b/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir
@@ -11,7 +11,7 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 1: index, 1: index] } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} { // CHECK-LABEL: spv.module // CHECK: spv.GlobalVariable @__push_constant_var__ : !spv.ptr<!spv.struct<(!spv.array<5 x i32, stride=4> [0])>, PushConstant> // CHECK: spv.func @push_constant() @@ -47,7 +47,7 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 1: index, 1: index] } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} { // CHECK-LABEL: spv.module // CHECK: spv.GlobalVariable @[[ARG0:.+]] bind(1, 2) : !spv.ptr<!spv.struct<(!spv.array<16 x f32, stride=4> [0])>, StorageBuffer> // CHECK: spv.GlobalVariable @[[ARG1_0:.+]] bind(1, 3) {aliased} : !spv.ptr<!spv.struct<(!spv.array<16 x f32, stride=4> [0])>, StorageBuffer> @@ -109,7 +109,7 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 1: index, 1: index] } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} { // CHECK-LABEL: spv.module // CHECK: spv.GlobalVariable @[[FUNC1_ARG:.+]] bind(1, 2) : !spv.ptr<!spv.struct<(!spv.array<16 x f32, stride=4> [0])>, StorageBuffer> // CHECK: spv.GlobalVariable @[[FUNC1_RET:.+]] bind(3, 4) : !spv.ptr<!spv.struct<(!spv.array<4 x vector<4xf32>, stride=16> [0])>, StorageBuffer> @@ -166,7 +166,7 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 1: index, 1: index] } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, SwiftShader:CPU, {}>} { func @interface_binding() { %c0 = constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<8x5xf32> @@ -213,7 +213,7 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 1: index, 1: index] } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, SwiftShader:CPU, {}>} { func @interface_wg_id() { %0 = hal.interface.workgroup.id[0] : index %1 = hal.interface.workgroup.id[1] : index @@ -252,7 +252,7 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 1: index, 1: index] } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, SwiftShader:CPU, {}>} { func @interface_wg_count() { %0 = hal.interface.workgroup.count[0] : index %1 = hal.interface.workgroup.count[1] : index
diff --git a/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir b/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir index 8ba2fd0..b8576ab 100644 --- a/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir +++ b/iree/compiler/Codegen/SPIRV/test/distribute_to_global_id.mlir
@@ -12,7 +12,7 @@ interface = @io, ordinal = 0 : index } - module attributes { + builtin.module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, @@ -88,7 +88,7 @@ interface = @io, ordinal = 0 : index } - module attributes { + builtin.module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, @@ -160,7 +160,7 @@ interface = @io, ordinal = 0 : index } - module attributes { + builtin.module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, @@ -209,7 +209,7 @@ interface = @io, ordinal = 0 : index } - module { + builtin.module { func @reduce_sum() { %c0 = constant 0 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<40x50x75xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir b/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir index 9756ed9..9bba3e7 100644 --- a/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir +++ b/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir
@@ -14,7 +14,7 @@ %z = constant 1: index hal.return %x, %y, %z: index, index, index } - module { + builtin.module { func @fold_block_id() -> (index, index, index) { %0 = "gpu.block_id"() {dimension = "x"} : () -> index %1 = "gpu.block_id"() {dimension = "y"} : () -> index @@ -49,7 +49,7 @@ %z = constant 1: index hal.return %x, %y, %z: index, index, index } - module { + builtin.module { func @fold_interface_workgroup_id() -> (index, index, index) { %0 = hal.interface.workgroup.id[0] : index %1 = hal.interface.workgroup.id[1] : index @@ -79,7 +79,7 @@ ordinal = 0 : index, workgroup_size = [8: index, 2: index, 1: index] } - module { + builtin.module { func @fold_thread_id() -> (index, index, index) { %0 = "gpu.thread_id"() {dimension = "x"} : () -> index %1 = "gpu.thread_id"() {dimension = "y"} : () -> index @@ -109,7 +109,7 @@ ordinal = 0 : index, workgroup_size = [8: index, 2: index, 1: index] } - module { + builtin.module { func @does_not_fold_mod() -> index { %0 = "gpu.thread_id"() {dimension = "z"} : () -> index %1 = affine.min affine_map<()[s0] -> (21, s0 mod 5)>()[%0] @@ -132,7 +132,7 @@ ordinal = 0 : index, workgroup_size = [8: index, 2: index, 1: index] } - module { + builtin.module { func @does_not_fold_div() -> index { %0 = "gpu.thread_id"() {dimension = "z"} : () -> index %1 = affine.min affine_map<()[s0] -> (21, s0 ceildiv 5)>()[%0] @@ -155,7 +155,7 @@ ordinal = 0 : index, workgroup_size = [8: index, 2: index, 1: index] } - module { + builtin.module { func @does_not_fold_symbol_mul_symbol() -> index { // 5 is in %0's range of [0,7] so we cannot fold the following into 5 or 0. %0 = "gpu.thread_id"() {dimension = "z"} : () -> index
diff --git a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir index 71ea9e1..045ce99 100644 --- a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir +++ b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir
@@ -29,7 +29,7 @@ interface = @io, ordinal = 0 : index } - module { + builtin.module { func @matmul_cooperative_matrix() { %c32 = constant 32 : index %c4096 = constant 4096 : index
diff --git a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir index ec3fcbf..dd4d216 100644 --- a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir +++ b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir
@@ -15,7 +15,7 @@ workgroup_size = [16: index, 1: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [64, 8]} } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, ARM:IntegratedGPU, {}>} { func @fuse_and_vectorize_fill_matmul() { %c0 = constant 0 : index %cst = constant 0.000000e+00 : f32 @@ -85,7 +85,7 @@ workgroup_size = [16: index, 1: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [64, 8]} } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, ARM:IntegratedGPU, {}>} { func @fuse_and_vectorize_matmul_add() { %c0 = constant 0 : index %cst = constant 0.000000e+00 : f32
diff --git a/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir b/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir index f996f1f..f15772d 100644 --- a/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir +++ b/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir
@@ -12,7 +12,7 @@ interface = @io, ordinal = 0 : index, workgroup_size = [16: index, 8: index, 1: index] } - module attributes { + builtin.module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 128 : i32, @@ -90,7 +90,7 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 4: index, 1: index] } - module attributes { + builtin.module attributes { spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 128 : i32,
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir index 6c1529b..67c2791 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
@@ -26,7 +26,11 @@ workgroup_size = [16: index, 8: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [8, 16]} } - module { + builtin.module attributes { + spv.target_env = + #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, + {max_compute_workgroup_invocations = 128 : i32, + max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} { func @matmul() { %c0 = constant 0 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<?x?xf32> @@ -105,7 +109,7 @@ workgroup_size = [32: index, 4: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [32, 4, 1]} } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative], [SPV_KHR_storage_buffer_storage_class]>, SwiftShader:CPU, {cooperative_matrix_properties_nv = [], max_compute_shared_memory_size = 16384 : i32, max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>, subgroup_size = 4 : i32}>} { func @conv_1d() { %cst = constant 0.000000e+00 : f32 %c0 = constant 0 : index @@ -192,7 +196,11 @@ workgroup_size = [32: index, 4: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [32, 4, 1]} } - module { + builtin.module attributes { + spv.target_env = + #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, + {max_compute_workgroup_invocations = 128 : i32, + max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} { func @conv_no_padding() { %c0 = constant 0 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<?x?x?x?xf32> @@ -317,7 +325,7 @@ workgroup_size = [32: index, 4: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [32, 4, 1]} } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader, GroupNonUniform, GroupNonUniformVote, GroupNonUniformArithmetic, GroupNonUniformBallot, GroupNonUniformShuffle, GroupNonUniformShuffleRelative], [SPV_KHR_storage_buffer_storage_class]>, SwiftShader:CPU, {cooperative_matrix_properties_nv = [], max_compute_shared_memory_size = 16384 : i32, max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>, subgroup_size = 4 : i32}>} { func @conv_3d() { %cst = constant 0.000000e+00 : f32 %c0 = constant 0 : index @@ -394,7 +402,7 @@ workgroup_size = [32: index, 4: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [32, 4, 1]} } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} { func @pooling_nhwc_max() { %c0 = constant 0 : index %0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<2x16x16x6xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir index 73fbce0..d17b02a 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
@@ -15,7 +15,7 @@ workgroup_size = [16: index, 1: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [64, 8, 1]} } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, ARM:IntegratedGPU, {}>} { func @batch_matmul_static_shape() { %c0 = constant 0 : index %c4 = constant 4 : index @@ -384,7 +384,7 @@ workgroup_size = [16: index, 1: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [64, 8, 1]} } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, ARM:IntegratedGPU, {}>} { func @fused_fill_batch_matmul() { %zero = constant 0.0 : f32 %c0 = constant 0 : index
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir index 2f72dd0..3fa302b 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
@@ -22,7 +22,7 @@ %z = constant 28: index hal.return %x, %y, %z: index, index, index } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, ARM:IntegratedGPU, {}>} { func @conv_static_shape_f32() { %cst = constant 0.000000e+00 : f32 %c32 = constant 32 : index @@ -122,7 +122,7 @@ %z = constant 14: index hal.return %x, %y, %z: index, index, index } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, ARM:IntegratedGPU, {}>} { func @depthwise_conv_static_shape_f32() { %cst = constant 0.000000e+00 : f32 %c96 = constant 96 : index
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir index 94c32fc..545757a 100644 --- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
@@ -15,7 +15,7 @@ workgroup_size = [16: index, 1: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [64, 8]} } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, ARM:IntegratedGPU, {}>} { func @matmul_static_shape_f16() { %cst = constant 0.000000e+00 : f16 %c0 = constant 0 : index @@ -82,7 +82,7 @@ workgroup_size = [16: index, 1: index, 1: index], translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [64, 8]} } - module { + builtin.module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], [SPV_KHR_storage_buffer_storage_class]>, ARM:IntegratedGPU, {}>} { func @matmul_static_shape_f32() { %c0 = constant 0 : index %cst = constant 0.000000e+00 : f32
diff --git a/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir b/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir index 97f1546..922d627 100644 --- a/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir +++ b/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir
@@ -12,7 +12,7 @@ [Shader, CooperativeMatrixNV, Int8, StorageBuffer8BitAccess], [SPV_KHR_storage_buffer_storage_class, SPV_NV_cooperative_matrix, SPV_KHR_8bit_storage]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>}> { - module { + builtin.module { // CHECK-LABEL: func @kernel_matmul func @kernel_matmul(%arg0: memref<8x32xi8>, %arg1: memref<32x8xi8>, %arg2: memref<8x8xi32>) attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { %c0 = constant 0 : index @@ -48,7 +48,7 @@ [Shader, CooperativeMatrixNV, Int8, Float16, StorageUniform16, StorageBuffer8BitAccess, Float16Buffer], [SPV_KHR_storage_buffer_storage_class, SPV_NV_cooperative_matrix, SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>}> { - module { + builtin.module { // CHECK-LABEL: func @kernel_matmul_licm func @kernel_matmul_licm(%arg0: memref<4096x4096xi8>, %arg1: memref<4096x4096xi8>, %arg2: memref<4096x4096xi32>) attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { %c32 = constant 32 : index @@ -95,7 +95,7 @@ [Shader, CooperativeMatrixNV, Int8, Float16, StorageUniform16, StorageBuffer8BitAccess, Float16Buffer], [SPV_KHR_storage_buffer_storage_class, SPV_NV_cooperative_matrix, SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>}> { - module { + builtin.module { // CHECK-LABEL: func @kernel_matmul_vector_memref func @kernel_matmul_vector_memref(%arg0: memref<4096x256xvector<4xi32>>, %arg1: memref<4096x256xvector<4xi32>>, %arg2: memref<4096x1024xvector<4xi32>>) attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { %c32 = constant 32 : index @@ -120,4 +120,4 @@ } } } -} \ No newline at end of file +}
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir index abef3ba..4f3f2fe 100644 --- a/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir +++ b/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
@@ -19,7 +19,10 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 1: index, 1: index] } - module { + builtin.module attributes { + spv.target_env = + #spv.target_env<#spv.vce<v1.5, [Shader], []>, + NVIDIA:DiscreteGPU, {subgroup_size = 32 : i32}>} { func @elementwise_static_shape() { %c0 = constant 0 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<128xf32> @@ -71,7 +74,10 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 1: index, 1: index] } - module { + builtin.module attributes { + spv.target_env = + #spv.target_env<#spv.vce<v1.5, [Shader], []>, + NVIDIA:DiscreteGPU, {subgroup_size = 32 : i32}>} { func @elementwise_transpose() { %c0 = constant 0 : index %arg0 = hal.interface.binding.subspan @io::@arg0[%c0] : memref<128x8xf32>
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir index d62ea6d..f7ae29c 100644 --- a/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir +++ b/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
@@ -37,7 +37,7 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 1: index, 1: index] } - module { + builtin.module { func @matmul_static_shape() { %c32 = constant 32 : index %c4096 = constant 4096 : index @@ -299,7 +299,7 @@ interface = @io, ordinal = 0 : index, workgroup_size = [32: index, 1: index, 1: index] } - module { + builtin.module { func @matmul_static_shape() { %c32 = constant 32 : index %c4096 = constant 4096 : index
diff --git a/iree/compiler/Dialect/Flow/IR/FlowOps.cpp b/iree/compiler/Dialect/Flow/IR/FlowOps.cpp index 189e717..6513688 100644 --- a/iree/compiler/Dialect/Flow/IR/FlowOps.cpp +++ b/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
@@ -553,7 +553,7 @@ } static void printExecutableOp(OpAsmPrinter &p, ExecutableOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.sym_name()); p.printOptionalAttrDictWithKeyword( op->getAttrs(), @@ -600,7 +600,7 @@ } static void printDispatchEntryOp(OpAsmPrinter &p, DispatchEntryOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.function_ref()); if (op.sym_name() != op.function_ref()) { p << " as(\"" << op.sym_name() << "\")";
diff --git a/iree/compiler/Dialect/Flow/IR/FlowOps.td b/iree/compiler/Dialect/Flow/IR/FlowOps.td index 09e1dfc..2df9623 100644 --- a/iree/compiler/Dialect/Flow/IR/FlowOps.td +++ b/iree/compiler/Dialect/Flow/IR/FlowOps.td
@@ -129,7 +129,7 @@ } def FLOW_DispatchWorkgroupRankOp : FLOW_PureOp<"dispatch.workgroup.rank", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{returns the rank of the workgroup dimensions}]; let description = [{ @@ -150,7 +150,7 @@ } def FLOW_DispatchWorkgroupIDOp : FLOW_PureOp<"dispatch.workgroup.id", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{returns the index of the current workgroup in the grid}]; let description = [{ @@ -182,7 +182,7 @@ } def FLOW_DispatchWorkgroupCountOp : FLOW_PureOp<"dispatch.workgroup.count", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{returns the total workgroup count of the grid}]; let description = [{ @@ -213,7 +213,7 @@ } def FLOW_DispatchWorkgroupSizeOp : FLOW_PureOp<"dispatch.workgroup.size", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{returns the size of each workgroup in invocations}]; let description = [{ @@ -252,7 +252,7 @@ } def FLOW_DispatchShapeOp : FLOW_PureOp<"dispatch.shape", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, DeclareOpInterfaceMethods<InferTypeOpInterface> ]> { let summary = [{returns the shape of a dispatch region input/output tensor}];
diff --git a/iree/compiler/Dialect/Flow/IR/test/dispatch_ops.mlir b/iree/compiler/Dialect/Flow/IR/test/dispatch_ops.mlir index 298bb75..8250528 100644 --- a/iree/compiler/Dialect/Flow/IR/test/dispatch_ops.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/dispatch_ops.mlir
@@ -1,7 +1,7 @@ // RUN: iree-opt -split-input-file %s | iree-opt -split-input-file | IreeFileCheck %s flow.executable @ex0 { - module { + builtin.module { func @dispatch_fn(%cst : index, %arg0 : tensor<4xf32>) -> tensor<4xf32> { return %arg0 : tensor<4xf32> }
diff --git a/iree/compiler/Dialect/Flow/IR/test/executable_ops.mlir b/iree/compiler/Dialect/Flow/IR/test/executable_ops.mlir index 4494dff..3369a78 100644 --- a/iree/compiler/Dialect/Flow/IR/test/executable_ops.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/executable_ops.mlir
@@ -5,7 +5,7 @@ // CHECK-LABEL: @dispatch_ex flow.executable @dispatch_ex { // CHECK: module { - module { + builtin.module { // CHECK: @dispatch0 func @dispatch0() { return
diff --git a/iree/compiler/Dialect/Flow/IR/test/stream_ops.mlir b/iree/compiler/Dialect/Flow/IR/test/stream_ops.mlir index f7323f5..89f69bc 100644 --- a/iree/compiler/Dialect/Flow/IR/test/stream_ops.mlir +++ b/iree/compiler/Dialect/Flow/IR/test/stream_ops.mlir
@@ -2,7 +2,7 @@ flow.executable @dispatch_0 { flow.dispatch.entry @rgn_dispatch_0 - module { + builtin.module { func @rgn_dispatch_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.multiply %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32>
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir b/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir index d5eaee6..37d9878 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/deduplicate_executables.mlir
@@ -3,7 +3,7 @@ // CHECK-LABEL: flow.executable @single_executable_ex_0 flow.executable @single_executable_ex_0 { flow.dispatch.entry @single_executable_entry_0 - module { + builtin.module { func @single_executable_entry_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.add %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32> @@ -23,7 +23,7 @@ // CHECK-LABEL: flow.executable @duplicate_executables_ex_0 flow.executable @duplicate_executables_ex_0 { flow.dispatch.entry @duplicate_executables_entry_0 - module { + builtin.module { func @duplicate_executables_entry_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.add %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32> @@ -33,7 +33,7 @@ // CHECK-NOT: flow.executable @duplicate_executables_ex_1 flow.executable @duplicate_executables_ex_1 { flow.dispatch.entry @duplicate_executables_entry_1 - module { + builtin.module { func @duplicate_executables_entry_1(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.add %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32> @@ -43,7 +43,7 @@ // CHECK-LABEL: flow.executable @duplicate_executables_ex_2 flow.executable @duplicate_executables_ex_2 { flow.dispatch.entry @duplicate_executables_entry_2 - module { + builtin.module { func @duplicate_executables_entry_2(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.subtract %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32> @@ -67,7 +67,7 @@ // CHECK: flow.executable @same_ops_diff_operands_ex_0 flow.executable @same_ops_diff_operands_ex_0 { flow.dispatch.entry @entry_0 - module { + builtin.module { func @entry_0(%arg0: tensor<2xi32>, %arg1: tensor<2xi32>) -> tensor<2xi32> { %0 = mhlo.multiply %arg0, %arg1 : tensor<2xi32> return %0 : tensor<2xi32> @@ -77,7 +77,7 @@ // CHECK: flow.executable @same_ops_diff_operands_ex_1 flow.executable @same_ops_diff_operands_ex_1 { flow.dispatch.entry @entry_1 - module { + builtin.module { func @entry_1(%arg0: tensor<2xi32>) -> tensor<2xi32> { %0 = mhlo.multiply %arg0, %arg0 : tensor<2xi32> return %0 : tensor<2xi32> @@ -100,7 +100,7 @@ flow.executable @multiple_entry_points_ex_0 { flow.dispatch.entry @multiple_entry_points_0_entry_0 flow.dispatch.entry @multiple_entry_points_0_entry_1 - module { + builtin.module { func @multiple_entry_points_0_entry_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.add %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32> @@ -115,7 +115,7 @@ flow.executable @multiple_entry_points_ex_1 { flow.dispatch.entry @multiple_entry_points_1_entry_0 flow.dispatch.entry @multiple_entry_points_1_entry_1 - module { + builtin.module { func @multiple_entry_points_1_entry_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.add %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32> @@ -145,7 +145,7 @@ // CHECK-LABEL: flow.executable @different_types_float_ex flow.executable @different_types_float_ex { flow.dispatch.entry @different_types_float_entry - module { + builtin.module { func @different_types_float_entry(%arg0: tensor<4xf32>) -> tensor<4xi1> { %0 = "mhlo.compare"(%arg0, %arg0) {comparison_direction = "EQ"} : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xi1> return %0 : tensor<4xi1> @@ -155,7 +155,7 @@ // CHECK-LABEL: flow.executable @different_types_int_ex flow.executable @different_types_int_ex { flow.dispatch.entry @different_types_int_entry - module { + builtin.module { func @different_types_int_entry(%arg0: tensor<4xi32>) -> tensor<4xi1> { %0 = "mhlo.compare"(%arg0, %arg0) {comparison_direction = "EQ"} : (tensor<4xi32>, tensor<4xi32>) -> tensor<4xi1> return %0 : tensor<4xi1> @@ -177,7 +177,7 @@ // CHECK-LABEL: flow.executable @nested_ops_ex_0 flow.executable @nested_ops_ex_0 { flow.dispatch.entry @nested_ops_entry_0 - module { + builtin.module { func @nested_ops_entry_0(%input: tensor<1x4xi32>) -> tensor<1xi32> { %0 = constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { @@ -192,7 +192,7 @@ // CHECK-NOT: flow.executable @nested_ops_ex_1 flow.executable @nested_ops_ex_1 { flow.dispatch.entry @nested_ops_entry_1 - module { + builtin.module { func @nested_ops_entry_1(%input: tensor<1x4xi32>) -> tensor<1xi32> { %0 = constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { @@ -207,7 +207,7 @@ // CHECK-LABEL: flow.executable @nested_ops_ex_2 flow.executable @nested_ops_ex_2 { flow.dispatch.entry @nested_ops_entry_2 - module { + builtin.module { func @nested_ops_entry_2(%input: tensor<1x4xi32>) -> tensor<1xi32> { %0 = constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { @@ -236,7 +236,7 @@ // CHECK-LABEL: flow.executable @attributes_ex_0 flow.executable @attributes_ex_0 { flow.dispatch.entry @attributes_entry_0 - module { + builtin.module { func @attributes_entry_0(%input: tensor<1x4xi32>) -> tensor<1xi32> { %0 = constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { @@ -252,7 +252,7 @@ // CHECK-LABEL: flow.executable @attributes_ex_1 flow.executable @attributes_ex_1 { flow.dispatch.entry @attributes_entry_1 - module { + builtin.module { func @attributes_entry_1(%input: tensor<1x4xi32>) -> tensor<1xi32> { %0 = constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { @@ -269,7 +269,7 @@ // CHECK-NOT: flow.executable @attributes_ex_2 flow.executable @attributes_ex_2 { flow.dispatch.entry @attributes_entry_2 - module { + builtin.module { func @attributes_entry_2(%input: tensor<1x4xi32>) -> tensor<1xi32> { %0 = constant dense<0> : tensor<i32> %1 = "mhlo.reduce"(%input, %0) ( { @@ -287,7 +287,7 @@ // CHECK-LABEL: flow.executable @block_successors_ex_0 flow.executable @block_successors_ex_0 { flow.dispatch.entry @entry_0 - module { + builtin.module { func @entry_0(%arg0: i32, %arg1: i32) -> i32 { %c0 = constant 0 : i32 %c1 = constant 1 : i32 @@ -303,7 +303,7 @@ // CHECK-LABEL: flow.executable @block_successors_ex_with_swapped_cond_br flow.executable @block_successors_ex_with_swapped_cond_br { flow.dispatch.entry @entry_1 - module { + builtin.module { func @entry_0(%arg0: i32, %arg1: i32) -> i32 { %c0 = constant 0 : i32 %c1 = constant 1 : i32
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/form_streams.mlir b/iree/compiler/Dialect/Flow/Transforms/test/form_streams.mlir index 069bd70..7e73103 100644 --- a/iree/compiler/Dialect/Flow/Transforms/test/form_streams.mlir +++ b/iree/compiler/Dialect/Flow/Transforms/test/form_streams.mlir
@@ -22,7 +22,7 @@ flow.dispatch.entry @outerOps_rgn_dispatch_0 attributes { workload = 4 : index } - module { + builtin.module { func @outerOps_rgn_dispatch_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.add %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32> @@ -75,7 +75,7 @@ flow.dispatch.entry @interleavedOuterOps_rgn_dispatch_0 attributes { workload = 4 : index } - module { + builtin.module { func @interleavedOuterOps_rgn_dispatch_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.add %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32> @@ -137,7 +137,7 @@ flow.dispatch.entry @interleavedDot_rgn_dispatch_0 attributes { workload = 16 : index } - module { + builtin.module { func @interleavedDot_rgn_dispatch_0(%arg0: tensor<4x4xf32>) -> tensor<4x4xf32> { %0 = mhlo.add %arg0, %arg0 : tensor<4x4xf32> return %0 : tensor<4x4xf32> @@ -148,7 +148,7 @@ flow.dispatch.entry @interleavedDot_rgn_dispatch_1 attributes { workload = 16 : index } - module { + builtin.module { func @interleavedDot_rgn_dispatch_1(%arg0: tensor<4x4xf32>, %arg1: tensor<4x4xf32>) -> tensor<4x4xf32> { %0 = "mhlo.dot"(%arg0, %arg1) : (tensor<4x4xf32>, tensor<4x4xf32>) -> tensor<4x4xf32> return %0 : tensor<4x4xf32> @@ -159,7 +159,7 @@ flow.dispatch.entry @interleavedDot_rgn_dispatch_2 attributes { workload = 16 : index } - module { + builtin.module { func @interleavedDot_rgn_dispatch_2(%arg0: tensor<4x4xf32>, %arg1: tensor<4x4xf32>) -> tensor<4x4xf32> { %0 = mhlo.multiply %arg0, %arg1 : tensor<4x4xf32> return %0 : tensor<4x4xf32> @@ -190,7 +190,7 @@ flow.dispatch.entry @caller_rgn_dispatch_0 attributes { workload = 4 : index } - module { + builtin.module { func @caller_rgn_dispatch_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.add %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32> @@ -201,7 +201,7 @@ flow.dispatch.entry @caller_rgn_dispatch_1 attributes { workload = 4 : index } - module { + builtin.module { func @caller_rgn_dispatch_1(%arg0: tensor<4xf32>, %arg1: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.multiply %arg0, %arg1 : tensor<4xf32> return %0 : tensor<4xf32> @@ -232,7 +232,7 @@ } flow.executable @callee_ex_dispatch_0 { flow.dispatch.entry @callee_rgn_dispatch_0 - module { + builtin.module { func @callee_rgn_dispatch_0(%arg0: tensor<4xf32>) -> tensor<4xf32> { %0 = mhlo.multiply %arg0, %arg0 : tensor<4xf32> return %0 : tensor<4xf32>
diff --git a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir index 55c7034..eed33d2 100644 --- a/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir +++ b/iree/compiler/Dialect/HAL/Conversion/FlowToHAL/test/stream_ops.mlir
@@ -12,7 +12,7 @@ interface = @interface, ordinal = 0 : index } - module {} + builtin.module {} } } @@ -172,7 +172,7 @@ interface = @interface, ordinal = 0 : index } - module {} + builtin.module {} } } @@ -308,7 +308,7 @@ interface = @interface, ordinal = 0 : index } - module {} + builtin.module {} } } @@ -371,7 +371,7 @@ interface = @io, ordinal = 0 : index } - module {} + builtin.module {} } } @@ -419,7 +419,7 @@ interface = @io, ordinal = 0 : index } - module {} + builtin.module {} } } @@ -479,7 +479,7 @@ interface = @interface_io, ordinal = 0 : index } - module {} + builtin.module {} } } @@ -493,7 +493,7 @@ interface = @interface_io, ordinal = 0 : index } - module {} + builtin.module {} } }
diff --git a/iree/compiler/Dialect/HAL/IR/HALOps.cpp b/iree/compiler/Dialect/HAL/IR/HALOps.cpp index e5f26b8..c1306dd 100644 --- a/iree/compiler/Dialect/HAL/IR/HALOps.cpp +++ b/iree/compiler/Dialect/HAL/IR/HALOps.cpp
@@ -535,7 +535,7 @@ static void printCommandBufferExecutionBarrierOp( OpAsmPrinter &p, CommandBufferExecutionBarrierOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printOperand(op.command_buffer()); p << ", \""; p << stringifyExecutionStageBitfield(op.source_stage_mask()); @@ -621,7 +621,7 @@ } static void printConstantPoolOp(OpAsmPrinter &p, ConstantPoolOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.sym_name()); p.printOptionalAttrDictWithKeyword( op->getAttrs(), @@ -782,7 +782,7 @@ } static void printDeviceSwitchOp(OpAsmPrinter &p, DeviceSwitchOp op) { - p << op.getOperationName() << "<"; + p << "<"; p.printOperand(op.device()); p << " : "; p.printType(op.device().getType()); @@ -863,7 +863,7 @@ } static void printExecutableOp(OpAsmPrinter &p, ExecutableOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.sym_name()); p.printOptionalAttrDictWithKeyword( op->getAttrs(), @@ -905,7 +905,7 @@ static void printExecutableEntryPointOp(OpAsmPrinter &p, ExecutableEntryPointOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.sym_name()); p.printOptionalAttrDictWithKeyword(op->getAttrs(), /*elidedAttrs=*/{"sym_name"}); @@ -983,7 +983,7 @@ } static void printExecutableVariantOp(OpAsmPrinter &p, ExecutableVariantOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.sym_name()); p << ", target = " << op.target(); p.printOptionalAttrDictWithKeyword( @@ -1034,7 +1034,7 @@ } static void printExecutableBinaryOp(OpAsmPrinter &p, ExecutableBinaryOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.sym_name()); p.printOptionalAttrDictWithKeyword( op->getAttrs(), @@ -1095,7 +1095,7 @@ } static void printInterfaceOp(OpAsmPrinter &p, InterfaceOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.sym_name()); p.printOptionalAttrDictWithKeyword( op->getAttrs(), @@ -1182,7 +1182,7 @@ } static void printInterfaceBindingOp(OpAsmPrinter &p, InterfaceBindingOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.sym_name()); p << ", set=" << op.set(); p << ", binding=" << op.binding();
diff --git a/iree/compiler/Dialect/HAL/IR/HALOps.td b/iree/compiler/Dialect/HAL/IR/HALOps.td index f93188d..ced6626 100644 --- a/iree/compiler/Dialect/HAL/IR/HALOps.td +++ b/iree/compiler/Dialect/HAL/IR/HALOps.td
@@ -29,7 +29,7 @@ // TODO(benvanik): remove these as the sequencer/other HAL ops are added. def HAL_ExSharedDeviceOp : HAL_PureOp<"ex.shared_device", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let results = (outs HAL_Device:$result @@ -113,7 +113,7 @@ //===----------------------------------------------------------------------===// def HAL_AllocatorComputeSizeOp : HAL_PureOp<"allocator.compute_size", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{buffer allocation size computation operation}]; let description = [{ @@ -160,7 +160,7 @@ } def HAL_AllocatorComputeOffsetOp : HAL_PureOp<"allocator.compute_offset", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, SameVariadicOperandSize, ]> { let summary = [{buffer view indices to byte offset computation operation}]; @@ -211,7 +211,7 @@ } def HAL_AllocatorComputeRangeOp : HAL_PureOp<"allocator.compute_range", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, SameVariadicOperandSize, ]> { let summary = [{buffer view byte range computation operation}]; @@ -267,7 +267,7 @@ } def HAL_AllocatorAllocateOp : HAL_Op<"allocator.allocate", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, DeclareOpInterfaceMethods<Util_SizeAwareOp>, ]> { let summary = [{empty buffer allocation operation}]; @@ -298,7 +298,7 @@ } def HAL_AllocatorConstantOp : HAL_Op<"allocator.constant", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{constant buffer allocation operation}]; let description = [{ @@ -331,7 +331,7 @@ } def HAL_AllocatorMapOp : HAL_Op<"allocator.map", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, DeclareOpInterfaceMethods<Util_SizeAwareOp>, ]> { let summary = [{allocator-supported host buffer wrapping operation}]; @@ -366,7 +366,7 @@ } def HAL_AllocatorPackOp : HAL_PureOp<"allocator.pack", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, AttrSizedOperandSegments, ]> { let summary = [{packs variable-sized slices into a single slab}]; @@ -474,7 +474,7 @@ //===----------------------------------------------------------------------===// def HAL_BufferAllocatorOp : HAL_PureOp<"buffer.allocator", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{buffer allocator accessor operation}]; let description = [{ @@ -498,7 +498,7 @@ } def HAL_BufferSubspanOp : HAL_PureOp<"buffer.subspan", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, DeclareOpInterfaceMethods<Util_SizeAwareOp>, ]> { let summary = [{buffer subspan operation}]; @@ -524,7 +524,7 @@ } def HAL_BufferLengthOp : HAL_PureOp<"buffer.length", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{buffer byte length accessor}]; let description = [{ @@ -594,7 +594,7 @@ //===----------------------------------------------------------------------===// def HAL_BufferViewCreateOp : HAL_PureOp<"buffer_view.create", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{buffer view reference initializer}]; let description = [{ @@ -641,7 +641,7 @@ } def HAL_BufferViewBufferOp : HAL_PureOp<"buffer_view.buffer", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{buffer view buffer accessor}]; let description = [{ @@ -665,7 +665,7 @@ } def HAL_BufferViewByteLengthOp : HAL_PureOp<"buffer_view.byte_length", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{buffer view buffer byte length accessor}]; let description = [{ @@ -822,7 +822,7 @@ //===----------------------------------------------------------------------===// def HAL_CommandBufferCreateOp : HAL_Op<"command_buffer.create", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{command buffer allocation operation}]; let description = [{ @@ -1346,7 +1346,7 @@ } def HAL_ConstantPoolLoadOp : HAL_PureOp<"constant_pool.load", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{constant pool tensor load pseudo-op}]; let description = [{ @@ -1399,7 +1399,7 @@ def HAL_ConstantStorageLookupOp : HAL_PureOp<"constant_storage.lookup", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{constant storage byte buffer accessor}]; let description = [{ @@ -1419,7 +1419,7 @@ } def HAL_ConstantSubspanOp : HAL_PureOp<"constant.subspan", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{runtime constant buffer lookup pseudo-op}]; let description = [{ @@ -1446,7 +1446,7 @@ //===----------------------------------------------------------------------===// def HAL_DescriptorSetCreateOp : HAL_PureOp<"descriptor_set.create", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, SameVariadicOperandSize, ]> { let summary = [{allocates a descriptor set from the device pool}]; @@ -1494,7 +1494,7 @@ def HAL_DescriptorSetLayoutCreateOp : HAL_PureOp<"descriptor_set_layout.create", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{creates a descriptor set layout}]; let description = [{ @@ -1523,7 +1523,7 @@ } def HAL_DescriptorSetLayoutLookupOp : HAL_PureOp<"descriptor_set_layout.lookup", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{descriptor set layout cache lookup pseudo-op}]; let description = [{ @@ -1554,7 +1554,7 @@ //===----------------------------------------------------------------------===// def HAL_DeviceAllocatorOp : HAL_PureOp<"device.allocator", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{device allocator accessor operation}]; let description = [{ @@ -1924,7 +1924,7 @@ } def HAL_ExecutableCreateOp : HAL_PureOp<"executable.create", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{creates an executable}]; let description = [{ @@ -1957,7 +1957,7 @@ } def HAL_ExecutableLookupOp : HAL_PureOp<"executable.lookup", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{executable cache lookup pseudo-op}]; let description = [{ @@ -2101,7 +2101,7 @@ } def HAL_InterfaceWorkgroupIDOp : HAL_PureOp<"interface.workgroup.id", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{returns the index of the current workgroup in the grid}]; let description = [{ @@ -2127,7 +2127,7 @@ } def HAL_InterfaceWorkgroupCountOp : HAL_PureOp<"interface.workgroup.count", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{returns the total workgroup count of the grid}]; let description = [{ @@ -2154,7 +2154,7 @@ } def HAL_InterfaceWorkgroupSizeOp : HAL_PureOp<"interface.workgroup.size", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{returns the size of each workgroup in invocations}]; let description = [{ @@ -2233,7 +2233,7 @@ //===----------------------------------------------------------------------===// def HAL_ExecutableLayoutCreateOp : HAL_PureOp<"executable_layout.create", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{creates an executable layout}]; let description = [{ @@ -2266,7 +2266,7 @@ } def HAL_ExecutableLayoutLookupOp : HAL_PureOp<"executable_layout.lookup", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{executable layout cache lookup pseudo-op}]; let description = [{ @@ -2299,7 +2299,7 @@ //===----------------------------------------------------------------------===// def HAL_SemaphoreCreateOp : HAL_Op<"semaphore.create", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{semaphore allocation operation}]; let description = [{
diff --git a/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir index b7e8fd9..fff3c6f 100644 --- a/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir
@@ -17,7 +17,7 @@ flow.dispatch.entry @add_dispatch_0 attributes { workgroup_rank = 3 : index } - module { + builtin.module { func @add_dispatch_0(%arg0: !flow.dispatch.tensor<readonly:16xf32>, %arg1: !flow.dispatch.tensor<readonly:16xf32>, %arg2: !flow.dispatch.tensor<writeonly:16xf32>) { %0 = linalg.init_tensor [16] : tensor<16xf32> %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir index de79b55..e6001f6 100644 --- a/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir
@@ -16,7 +16,7 @@ flow.dispatch.entry @add_dispatch_0 attributes { workgroup_rank = 3 : index } - module { + builtin.module { func @add_dispatch_0(%arg0: !flow.dispatch.tensor<readonly:16xf32>, %arg1: !flow.dispatch.tensor<readonly:16xf32>, %arg2: !flow.dispatch.tensor<writeonly:16xf32>) { %0 = linalg.init_tensor [16] : tensor<16xf32> %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir index 2bb125a..6628de1 100644 --- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir
@@ -18,7 +18,7 @@ flow.dispatch.entry @add_dispatch_0 attributes { workgroup_rank = 3 : index } - module { + builtin.module { func @add_dispatch_0(%arg0: !flow.dispatch.tensor<readonly:16xf32>, %arg1: !flow.dispatch.tensor<readonly:16xf32>, %arg2: !flow.dispatch.tensor<writeonly:16xf32>) { %0 = linalg.init_tensor [16] : tensor<16xf32> %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
diff --git a/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir index ab4185f..d0c2ae4 100644 --- a/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir
@@ -17,7 +17,7 @@ signature = (tensor<16xf32>, tensor<16xf32>) -> tensor<16xf32>, workgroup_rank = 3 : index } - module { + builtin.module { func @add_dispatch_0(%arg0: !flow.dispatch.tensor<readonly:16xf32>, %arg1: !flow.dispatch.tensor<readonly:16xf32>, %arg2: !flow.dispatch.tensor<writeonly:16xf32>) { %0 = linalg.init_tensor [16] : tensor<16xf32> %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir b/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir index dcc219c..7921db2 100644 --- a/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir +++ b/iree/compiler/Dialect/HAL/Target/VMVX/test/linking.mlir
@@ -10,7 +10,7 @@ } hal.executable.variant @vmvx, target = #vmvx_target { hal.executable.entry_point @dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { vm.module @module { vm.func @dispatch_0() { vm.return @@ -28,7 +28,7 @@ } hal.executable.variant @vmvx, target = #vmvx_target { hal.executable.entry_point @dispatch_1 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { vm.module @module { vm.func @dispatch_1() { vm.return @@ -47,7 +47,7 @@ } hal.executable.variant @vmvx, target = #vmvx_target { hal.executable.entry_point @dispatch_2 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { vm.module @module { vm.func @dispatch_2() { vm.return @@ -127,7 +127,7 @@ } hal.executable.variant @vmvx, target = #vmvx_target { hal.executable.entry_point @dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { vm.module @module { vm.func @dispatch_0() { vm.return @@ -137,7 +137,7 @@ } } hal.executable.variant @cuda, target = #cuda_target { - module { + builtin.module { } } } @@ -148,7 +148,7 @@ } hal.executable.variant @vmvx, target = #vmvx_target { hal.executable.entry_point @dispatch_1 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { vm.module @module { vm.func @dispatch_1() { vm.return @@ -158,7 +158,7 @@ } } hal.executable.variant @cuda, target = #cuda_target { - module { + builtin.module { } } } @@ -249,7 +249,7 @@ } hal.executable.variant @vmvx, target = #vmvx_target { hal.executable.entry_point @dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { vm.module @module {} } } @@ -262,7 +262,7 @@ } hal.executable.variant @vmvx, target = #vmvx_target { hal.executable.entry_point @dispatch_1 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { vm.module @module {} } } @@ -275,7 +275,7 @@ } hal.executable.variant @vmvx, target = #vmvx_target { hal.executable.entry_point @dispatch_2 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { vm.module @module {} } } @@ -306,7 +306,7 @@ hal.interface @io {} hal.executable.variant @vmvx, target = #vmvx_target { hal.executable.entry_point @dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { vm.module @module { vm.rodata public @rodata_a dense<[0]> : tensor<1xi32> vm.rodata public @rodata_b dense<[0]> : tensor<1xi32> @@ -329,7 +329,7 @@ hal.interface @io {} hal.executable.variant @vmvx, target = #vmvx_target { hal.executable.entry_point @dispatch_1 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { vm.module @module { // Conflict with a public symbol, this should be renamed when linked. vm.rodata private @rodata_b dense<[1]> : tensor<1xi32>
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir index f2e74bd..51fb789 100644 --- a/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
@@ -16,7 +16,7 @@ flow.dispatch.entry @entry attributes { workgroup_rank = 3 : index } - module { + builtin.module { func @entry(%arg0: !flow.dispatch.tensor<readonly:16xf32>, %arg1: !flow.dispatch.tensor<readonly:16xf32>, %arg2: !flow.dispatch.tensor<writeonly:16xf32>) { %0 = linalg.init_tensor [16] : tensor<16xf32> %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir index d333768..7ff3f3b 100644 --- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir +++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/linking.mlir
@@ -12,7 +12,7 @@ } hal.executable.variant @vulkan_spirv_fb, target = #executable_target_vulkan_spirv_fb { hal.executable.entry_point @call_dispatch_0 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> { spv.func @call_dispatch_0() "None" { spv.Return @@ -35,7 +35,7 @@ } hal.executable.variant @vulkan_spirv_fb, target = #executable_target_vulkan_spirv_fb { hal.executable.entry_point @call_dispatch_1 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> { spv.func @call_dispatch_1() "None" { spv.Return @@ -58,7 +58,7 @@ } hal.executable.variant @vulkan_spirv_fb, target = #executable_target_vulkan_spirv_fb { hal.executable.entry_point @call_dispatch_2 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> { spv.func @call_dispatch_2() "None" { spv.Return @@ -87,7 +87,7 @@ %c56_0 = constant 56 : index hal.return %c1, %c56, %c56_0 : index, index, index } - module { + builtin.module { spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> { spv.func @call_dispatch_3() "None" { spv.Return @@ -111,7 +111,7 @@ } hal.executable.variant @vulkan_spirv_fb, target = #executable_target_vulkan_spirv_fb { hal.executable.entry_point @call_dispatch_4 attributes {interface = @io, ordinal = 0 : index} - module { + builtin.module { spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> { spv.func @call_dispatch_4() "None" { spv.Return
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir index cbcb3dd..6590136 100644 --- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir +++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
@@ -18,7 +18,7 @@ flow.dispatch.entry @add_dispatch_0 attributes { workgroup_rank = 3 : index } - module { + builtin.module { func @add_dispatch_0(%arg0: !flow.dispatch.tensor<readonly:16xf32>, %arg1: !flow.dispatch.tensor<readonly:16xf32>, %arg2: !flow.dispatch.tensor<writeonly:16xf32>) { %0 = linalg.init_tensor [16] : tensor<16xf32> %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir b/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir index a2d64ab..e1961f2 100644 --- a/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir +++ b/iree/compiler/Dialect/HAL/Transforms/test/materialize_interfaces.mlir
@@ -19,7 +19,7 @@ workgroup_rank = 2 : index } // CHECK-NEXT: module { - module { + builtin.module { // CHECK-NEXT: func @entry() { func @entry(%arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<writeonly:4x8xf32>) { // CHECK-NEXT: %c0 = constant 0 : index @@ -79,7 +79,7 @@ workgroup_rank = 2 : index } // CHECK-NEXT: module { - module { + builtin.module { // CHECK-NEXT: func @entry() { func @entry(%arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<writeonly:4x8xf32>) { // CHECK-NEXT: %c0 = constant 0 : index @@ -138,7 +138,7 @@ workgroup_rank = 2 : index } // CHECK-NEXT: module { - module { + builtin.module { // CHECK-NEXT: func @entry() { func @entry( // CHECK-NEXT: %c0 = constant 0 : index @@ -217,7 +217,7 @@ workgroup_rank = 2 : index } // CHECK-NEXT: module { - module { + builtin.module { // CHECK-NEXT: func @entry() { func @entry(%arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<writeonly:4x8xf32>) { // CHECK-DAG: %[[WORKGROUP_ID_X:.+]] = hal.interface.workgroup.id[0] : index @@ -258,7 +258,7 @@ flow.dispatch.entry @entry attributes { workgroup_rank = 2 : index } - module { + builtin.module { // CHECK: func @entry() { func @entry(%arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<readwrite:4x8xf32>) { // CHECK-NEXT: %c0 = constant 0 : index @@ -312,7 +312,7 @@ flow.dispatch.entry @entry attributes { workgroup_rank = 2 : index } - module { + builtin.module { // CHECK: func @entry() { func @entry( %arg: !flow.dispatch.tensor<readonly:8x4xf32>,
diff --git a/iree/compiler/Dialect/HAL/Transforms/test/propagate_constant_workgroup_info.mlir b/iree/compiler/Dialect/HAL/Transforms/test/propagate_constant_workgroup_info.mlir index 3ed7a15..670542c 100644 --- a/iree/compiler/Dialect/HAL/Transforms/test/propagate_constant_workgroup_info.mlir +++ b/iree/compiler/Dialect/HAL/Transforms/test/propagate_constant_workgroup_info.mlir
@@ -11,7 +11,7 @@ ordinal = 0 : index, workgroup_size = [32 : index, 4 : index, 8 : index] } - module { + builtin.module { // CHECK: func @entry() func @entry() { // CHECK-DAG: constant 32 : index
diff --git a/iree/compiler/Dialect/Shape/IR/ShapeOps.cpp b/iree/compiler/Dialect/Shape/IR/ShapeOps.cpp index 806682c..ad41caa 100644 --- a/iree/compiler/Dialect/Shape/IR/ShapeOps.cpp +++ b/iree/compiler/Dialect/Shape/IR/ShapeOps.cpp
@@ -163,7 +163,7 @@ } static void printRankedDimOp(OpAsmPrinter &p, RankedDimOp op) { - p << op.getOperationName() << " "; + p << " "; p.printOperand(op.shape()); p << "[" << op.getIndex() << "]"; p << " : ";
diff --git a/iree/compiler/Dialect/Shape/IR/ShapeOps.td b/iree/compiler/Dialect/Shape/IR/ShapeOps.td index ca6c14c..a6b541d 100644 --- a/iree/compiler/Dialect/Shape/IR/ShapeOps.td +++ b/iree/compiler/Dialect/Shape/IR/ShapeOps.td
@@ -113,7 +113,7 @@ } def Shape_ConstRankedShapeOp : Shape_PureOp<"const_ranked_shape", - [ConstantLike, DeclareOpInterfaceMethods<OpAsmOpInterface>]> { + [ConstantLike, DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>]> { let summary = "A constant ranked_shape."; let description = [{ Holds a RankedShape value. Note that it is only legal to store a constant
diff --git a/iree/compiler/Dialect/Util/IR/UtilOps.cpp b/iree/compiler/Dialect/Util/IR/UtilOps.cpp index 4177d87..619cdde 100644 --- a/iree/compiler/Dialect/Util/IR/UtilOps.cpp +++ b/iree/compiler/Dialect/Util/IR/UtilOps.cpp
@@ -481,7 +481,6 @@ } void printDoNotOptimizeOp(OpAsmPrinter &p, Operation *op) { - p << "util.do_not_optimize"; p << "("; p.printOperands(op->getOperands()); p << ")"; @@ -538,7 +537,7 @@ void printUnfoldableConstantOp(OpAsmPrinter &p, Operation *op) { auto constOp = cast<IREE::Util::UnfoldableConstantOp>(op); - p << "util.unfoldable_constant "; + p << " "; p.printOptionalAttrDict(constOp->getAttrs(), /*elidedAttrs=*/{"value"}); if (constOp->getAttrs().size() > 1) p << ' '; @@ -595,7 +594,6 @@ } static void printInitializerOp(OpAsmPrinter &p, InitializerOp &op) { - p << "util.initializer"; p.printOptionalAttrDictWithKeyword(op->getAttrs(), /*elidedAttrs=*/{"type"}); p.printRegion(op.body()); }
diff --git a/iree/compiler/Dialect/Util/IR/UtilOps.td b/iree/compiler/Dialect/Util/IR/UtilOps.td index 131eef4..b188fe4 100644 --- a/iree/compiler/Dialect/Util/IR/UtilOps.td +++ b/iree/compiler/Dialect/Util/IR/UtilOps.td
@@ -278,7 +278,7 @@ } def Util_GlobalAddressOp : Util_PureOp<"global.address", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, ]> { let summary = [{returns an address reference to a global}]; let description = [{ @@ -303,7 +303,7 @@ } def Util_GlobalLoadOp : Util_Op<"global.load", [ - DeclareOpInterfaceMethods<OpAsmOpInterface>, + DeclareOpInterfaceMethods<OpAsmOpInterface, ["getAsmResultNames"]>, // HACK: works around the lack of symbol side effects in C++. DeclareOpInterfaceMethods<MemoryEffectsOpInterface>, ]> {
diff --git a/iree/compiler/Dialect/VM/IR/VMOps.cpp b/iree/compiler/Dialect/VM/IR/VMOps.cpp index 2bfb6c8..86b7eb4 100644 --- a/iree/compiler/Dialect/VM/IR/VMOps.cpp +++ b/iree/compiler/Dialect/VM/IR/VMOps.cpp
@@ -144,7 +144,7 @@ } static void printExportOp(OpAsmPrinter &p, ExportOp op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.function_ref()); if (op.export_name() != op.function_ref()) { p << " as(\"" << op.export_name() << "\")"; @@ -226,7 +226,7 @@ } static void printImportOp(OpAsmPrinter &p, ImportOp &op) { - p << op.getOperationName() << ' '; + p << ' '; p.printSymbolName(op.getName()); p << "("; for (int i = 0; i < op.getNumFuncArguments(); ++i) { @@ -304,7 +304,6 @@ } static void printInitializerOp(OpAsmPrinter &p, InitializerOp &op) { - p << "vm.initializer"; p.printOptionalAttrDictWithKeyword(op->getAttrs(), /*elidedAttrs=*/{"type"}); p.printRegion(op.body()); } @@ -479,7 +478,7 @@ template <typename T> static void printConstOp(OpAsmPrinter &p, T &op) { - p << op.getOperationName() << ' '; + p << ' '; p.printAttribute(op.value()); p.printOptionalAttrDict(op->getAttrs(), /*elidedAttrs=*/{"value"}); } @@ -814,7 +813,7 @@ template <typename T> static void printSwitchOp(OpAsmPrinter &p, T &op) { - p << op.getOperationName() << " "; + p << " "; p.printOperand(op.index()); p << "["; p.printOperands(op.values()); @@ -1022,7 +1021,7 @@ } static void printCallVariadicOp(OpAsmPrinter &p, CallVariadicOp &op) { - p << op.getOperationName() << ' ' << op->getAttr("callee") << '('; + p << ' ' << op->getAttr("callee") << '('; int operand = 0; llvm::interleaveComma( llvm::zip(op.segment_sizes(), op.segment_types()), p, @@ -1135,7 +1134,7 @@ } static void printCondFailOp(OpAsmPrinter &p, CondFailOp op) { - p << op.getOperationName() << ' '; + p << ' '; if (op.condition() != op.status()) { p << op.condition() << ", "; }
diff --git a/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir b/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir index aad52ed..4987e18 100644 --- a/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir +++ b/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
@@ -16,7 +16,7 @@ #map0 = affine_map<(d0) -> (d0)> flow.executable @simpleMath_dispatch_0 { flow.dispatch.entry @simpleMath_dispatch_0 attributes {workgroup_rank = 3 : index} - module { + builtin.module { func @simpleMath_dispatch_0(%arg0: !flow.dispatch.tensor<readonly:4xf32>, %arg1: !flow.dispatch.tensor<writeonly:4xf32>) { %c4 = constant 4 : index %c1 = constant 1 : index
diff --git a/iree/compiler/InputConversion/Common/test/iree_import_public.mlir b/iree/compiler/InputConversion/Common/test/iree_import_public.mlir index a8ea2ba..4cf1713 100644 --- a/iree/compiler/InputConversion/Common/test/iree_import_public.mlir +++ b/iree/compiler/InputConversion/Common/test/iree_import_public.mlir
@@ -192,7 +192,7 @@ // CHECK-NEXT: util.global.store %[[VALUE]], @global5 : tensor<4xi32> // CHECK-NEXT: util.initializer.return // CHECK-NEXT: } - // CHECK: builtin.func private @initializer() -> tensor<4xi32> + // CHECK: func private @initializer() -> tensor<4xi32> builtin.func private @initializer() -> tensor<4xi32> }
diff --git a/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir b/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir index 442c27b..38e602e 100644 --- a/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir +++ b/iree/compiler/InputConversion/MHLO/test/convert_mhlo_to_linalg_ext.mlir
@@ -317,7 +317,7 @@ // CHECK-SAME: outs(%[[R2]]#0, %[[R2]]#1 // CHECK: %[[RES_REAL:.+]] = tensor.extract_slice %[[R3]]#0[0] [5] [1] : tensor<8xf32> to tensor<5xf32> // CHECK: %[[RES_IMAG:.+]] = tensor.extract_slice %[[R3]]#1[0] [5] [1] : tensor<8xf32> to tensor<5xf32> -// CHECK: %{{.+}} = "mhlo.complex"(%[[RES_REAL]], %[[RES_IMAG]]) +// CHECK: %{{.+}} = mhlo.complex(%[[RES_REAL]], %[[RES_IMAG]]) // ----- @@ -366,4 +366,4 @@ // CHECK-SAME: outs(%[[R2]]#0, %[[R2]]#1 // CHECK: %[[RES_REAL:.+]] = tensor.extract_slice %[[R3]]#0[0, 0] [4, 5] [1, 1] : tensor<4x8xf32> to tensor<4x5xf32> // CHECK: %[[RES_IMAG:.+]] = tensor.extract_slice %[[R3]]#1[0, 0] [4, 5] [1, 1] : tensor<4x8xf32> to tensor<4x5xf32> -// CHECK: %{{.+}} = "mhlo.complex"(%[[RES_REAL]], %[[RES_IMAG]]) +// CHECK: %{{.+}} = mhlo.complex(%[[RES_REAL]], %[[RES_IMAG]])
diff --git a/third_party/llvm-project b/third_party/llvm-project index b9db703..9b6c813 160000 --- a/third_party/llvm-project +++ b/third_party/llvm-project
@@ -1 +1 @@ -Subproject commit b9db70369b7799887b817e13109801795e4d70fc +Subproject commit 9b6c8132d3785269512803ff51cb421f8d8bcf0e