[spirv] Add missing WIE patterns required for `--iree-mhlo-demote-i64-to-i32=false`

[kuhar]

When running the unittest suite with --iree-mhlo-demote-i64-to-i32=false, a few Wide Integer Emulation patterns appear missing: arith.maxsi, arith.minsi, arith.index_cast.

Full log:

$ ninja all iree-test-deps && ctest --output-on-failure -E vulkan_f16 -j64
[0/2] Re-checking globbed directories...[399/556] Generating check_vulkan-spirv_vulkan_gather.mlir_module.vmfb from gather.mlirFAILED: tests/e2e/xla_ops/check_vulkan-spirv_vulkan_gather.mlir_module.vmfb /usr/local/google/home/kubak/iree/build-relass/tests/e2e/xla_ops/check_vulkan-spirv_vulkan_gather.mlir_module.vmfb cd /usr/local/google/home/kubak/iree/build-relass/tests/e2e/xla_ops && /usr/local/google/home/kubak/iree/build-relass/tools/iree-compile --output-format=vm-bytecode --mlir-print-op-on-diagnostic=false --iree-hal-target-backends=vulkan-spirv --iree-input-type=mhlo /usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/gather.mlir -o check_vulkan-spirv_vulkan_gather.mlir_module.vmfb --iree-llvm-embedded-linker-path=\"/usr/local/google/home/kubak/iree/build-relass/third_party/llvm-project/llvm/bin/lld\" --iree-llvm-wasm-linker-path=\"/usr/local/google/home/kubak/iree/build-relass/third_party/llvm-project/llvm/bin/lld\"
/usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/gather.mlir:9:10: error: failed to legalize operation 'arith.index_cast' that was explicitly marked illegal
  %res = "mhlo.gather"(%input, %start_indices) {
         ^
/usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/gather.mlir:9:10: error: failed to run translation of source executable to target executable for backend #hal.executable.target<"vulkan", "vulkan-spirv-fb", {spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64, cooperative_matrix_properties_nv = []>>}>
  %res = "mhlo.gather"(%input, %start_indices) {
         ^/usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/gather.mlir:9:10: error: failed to serialize executables
  %res = "mhlo.gather"(%input, %start_indices) {
         ^
compilation failed
[418/556] Generating check_vulkan-spirv_vulkan_minimum.mlir_module.vmfb from minimum.mlir
FAILED: tests/e2e/xla_ops/check_vulkan-spirv_vulkan_minimum.mlir_module.vmfb /usr/local/google/home/kubak/iree/build-relass/tests/e2e/xla_ops/check_vulkan-spirv_vulkan_minimum.mlir_module.vmfb 
cd /usr/local/google/home/kubak/iree/build-relass/tests/e2e/xla_ops && /usr/local/google/home/kubak/iree/build-relass/tools/iree-compile --output-format=vm-bytecode --mlir-print-op-on-diagnostic=false --iree-hal-target-backends=vulkan-spirv --iree-input-type=mhlo /usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/minimum.mlir -o check_vulkan-spirv_vulkan_minimum.mlir_module.vmfb --iree-llvm-embedded-linker-path=\"/usr/local/google/home/kubak/iree/build-relass/third_party/llvm-project/llvm/bin/lld\" --iree-llvm-wasm-linker-path=\"/usr/local/google/home/kubak/iree/build-relass/third_party/llvm-project/llvm/bin/lld\"
/usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/minimum.mlir:52:13: error: failed to legalize operation 'arith.minsi' that was explicitly marked illegal
  %result = "mhlo.minimum"(%lhs, %rhs) : (tensor<i64>, tensor<i64>) -> tensor<i64>
            ^
/usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/minimum.mlir:52:13: error: failed to run translation of source executable to target executable for backend #hal.executable.target<"vulkan", "vulkan-spirv-fb", {spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64, cooperative_matrix_properties_nv = []>>}>
  %result = "mhlo.minimum"(%lhs, %rhs) : (tensor<i64>, tensor<i64>) -> tensor<i64>
            ^
/usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/minimum.mlir:52:13: error: failed to serialize executables
  %result = "mhlo.minimum"(%lhs, %rhs) : (tensor<i64>, tensor<i64>) -> tensor<i64>
            ^
compilation failed
[422/556] Generating check_vulkan-spirv_vulkan_maximum.mlir_module.vmfb from maximum.mlir
FAILED: tests/e2e/xla_ops/check_vulkan-spirv_vulkan_maximum.mlir_module.vmfb /usr/local/google/home/kubak/iree/build-relass/tests/e2e/xla_ops/check_vulkan-spirv_vulkan_maximum.mlir_module.vmfb 
cd /usr/local/google/home/kubak/iree/build-relass/tests/e2e/xla_ops && /usr/local/google/home/kubak/iree/build-relass/tools/iree-compile --output-format=vm-bytecode --mlir-print-op-on-diagnostic=false --iree-hal-target-backends=vulkan-spirv --iree-input-type=mhlo /usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/maximum.mlir -o check_vulkan-spirv_vulkan_maximum.mlir_module.vmfb --iree-llvm-embedded-linker-path=\"/usr/local/google/home/kubak/iree/build-relass/third_party/llvm-project/llvm/bin/lld\" --iree-llvm-wasm-linker-path=\"/usr/local/google/home/kubak/iree/build-relass/third_party/llvm-project/llvm/bin/lld\"
/usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/maximum.mlir:52:13: error: failed to legalize operation 'arith.maxsi' that was explicitly marked illegal
  %result = "mhlo.maximum"(%lhs, %rhs) : (tensor<i64>, tensor<i64>) -> tensor<i64>
            ^
/usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/maximum.mlir:52:13: error: failed to run translation of source executable to target executable for backend #hal.executable.target<"vulkan", "vulkan-spirv-fb", {spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Shader, GroupNonUniform], [SPV_KHR_storage_buffer_storage_class, SPV_KHR_variable_pointers]>, #spirv.resource_limits<max_compute_workgroup_size = [128, 128, 64], subgroup_size = 64, cooperative_matrix_properties_nv = []>>}>
  %result = "mhlo.maximum"(%lhs, %rhs) : (tensor<i64>, tensor<i64>) -> tensor<i64>
            ^
/usr/local/google/home/kubak/iree/iree/tests/e2e/xla_ops/maximum.mlir:52:13: error: failed to serialize executables
  %result = "mhlo.maximum"(%lhs, %rhs) : (tensor<i64>, tensor<i64>) -> tensor<i64>            ^
compilation failed
[528/556] Generating check_llvm-cpu_local-task_mobilenetv3_fake_weights.mlir_module.vmfb from mobilenetv3_fake_weights.mlir
ninja: build stopped: subcommand failed.

[kuhar]

cc: @antiagainst

[antiagainst]

Ouch.. Hopefully this is the last batch. Feel free to take on them when you want to slot something in; no hurry.

[kuhar]

I added the missing patterns in:

1. https://reviews.llvm.org/D138225
2. https://reviews.llvm.org/D138184

I'll wait for the next integrate to pick those up and see if everything builds with --iree-mhlo-demote-i64-to-i32=false

[kuhar]

I checked and the whole test suite builds with --iree-mhlo-demote-i64-to-i32=false with the most recent main branch.