<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=https://github.com/llvm/llvm-project/issues/74939>74939</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            [mlir::spirv] mlir-vulkan-runner fails on scf.if
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            mlir:spirv
      </td>
    </tr>

    <tr>
      <th>Assignees</th>
      <td>
      </td>
    </tr>

    <tr>
      <th>Reporter</th>
      <td>
          python3kgae
      </td>
    </tr>
</table>

<pre>
    lit fails when running following mlir-vulkan-runner test for scf.if.


```
// RUN: mlir-vulkan-runner %s --shared-libs=%vulkan-runtime-wrappers,%mlir_runner_utils --entry-point-result=void | FileCheck %s

// CHECK: [3.3,  3.3,  3.3, 3.3,  0,  0,  0,  0]
module attributes {
  gpu.container_module,
 spirv.target_env = #spirv.target_env<
    #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>
} {
 gpu.module @kernels {
    gpu.func @kernel_add(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
      kernel attributes { spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} {
      %0 = gpu.block_id x
      %limit = arith.constant 4 : index
 %cond = arith.cmpi slt, %0, %limit : index
      scf.if %cond {     
        %1 = memref.load %arg0[%0] : memref<8xf32>
        %2 = memref.load %arg1[%0] : memref<8xf32>
        %3 = arith.addf %1, %2 : f32
        memref.store %3, %arg2[%0] : memref<8xf32>
      }
 gpu.return
    }
  }

  func.func @main() {
    %arg0 = memref.alloc() : memref<8xf32>
    %arg1 = memref.alloc() : memref<8xf32>
    %arg2 = memref.alloc() : memref<8xf32>
    %0 = arith.constant 0 : i32
    %1 = arith.constant 1 : i32
    %2 = arith.constant 2 : i32
    %value0 = arith.constant 0.0 : f32
 %value1 = arith.constant 1.1 : f32
    %value2 = arith.constant 2.2 : f32
    %arg3 = memref.cast %arg0 : memref<8xf32> to memref<?xf32>
 %arg4 = memref.cast %arg1 : memref<8xf32> to memref<?xf32>
    %arg5 = memref.cast %arg2 : memref<8xf32> to memref<?xf32>
    call @fillResource1DFloat(%arg3, %value1) : (memref<?xf32>, f32) -> ()
    call @fillResource1DFloat(%arg4, %value2) : (memref<?xf32>, f32) -> ()
    call @fillResource1DFloat(%arg5, %value0) : (memref<?xf32>, f32) -> ()

 %cst1 = arith.constant 1 : index
    %cst8 = arith.constant 8 : index
 gpu.launch_func @kernels::@kernel_add
        blocks in (%cst8, %cst1, %cst1) threads in (%cst1, %cst1, %cst1)
        args(%arg0 : memref<8xf32>, %arg1 : memref<8xf32>, %arg2 : memref<8xf32>)
    %arg6 = memref.cast %arg5 : memref<?xf32> to memref<*xf32>
    call @printMemrefF32(%arg6) : (memref<*xf32>) -> ()
    return
  }
 func.func private @fillResource1DFloat(%0 : memref<?xf32>, %1 : f32)
[addf_if.mlir.txt](https://github.com/llvm/llvm-project/files/13624972/addf_if.mlir.txt)

 func.func private @printMemrefF32(%ptr : memref<*xf32>)
}
```

There are three different issues related to this problem.
1. mlir-vulkan-runner doesn't register SCFDialect,
2. convert SCFToSPIRV
   If put  createSCFToSPIRV after createConvertGPUToSPIRVPass. createConvertGPUToSPIRVPass will complain failed to legalize scf.if. 
   If put createSCFToSPIRV before createConvertGPUToSPIRVPass,  SCFToSPIRVPass will convert all func.func, so func.func 'main' in the test will be converted SPIRV too.
4. UnrealizedConversionCastOp cannot be removed when used in pattern like this.
```
    %5 = builtin.unrealized_conversion_cast %cst1_i32 : i32 to index
    %6 = builtin.unrealized_conversion_cast %5 : index to i64
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJy8WFtvozwa_jXOjRUEJiThIhdtMtlvNNrdqnO4RQZegreOjWyTzuyvX9kmHFLSUUfaL4pIwI-f9_zahmrNTgJgh5JHlBwWtDW1VLvml6mliF9OFBa5LH_tODO4ooxr_FqDwKoVgokTriTn8tX-O3OmlpeWv1CxtKOgsAFtcCUV1kUVsCpA4QGFD5PrOuy-_pYcETni5-__QvHDHCMiicbLpa6pgnLJWa5RfEAkGVCGnWH5qmjTgNKI7BFJLE_mCbLWWBOWSxBG_Vo2kgmzVKBbblB8uEhWYrTZ4yPjsK-heHECJwp7Dfd_fdp_sTqi5DEOYkT2GN_8Xm_DmWty8GxnWbYcMDVGsbw1oDHaPPohjE9NGxRSGMqs4h5qDfLDumHqEhiqTmAyEBeM4gNGJL59juLrDIyH8UsBKN5fosBphJLHrzUtQVnFuvunH9mXv54zbaSiJ8jytqpA9bcFp1pbdPzJTbjyKtCyVQVknJ2Z0VZ4_Ml-vfM2h5GB1r7OAWgVvoASwCcO8C6oWlEMgIyWJSJbRBKqTiF2aQJnBRWK99ufVUx6jSwg-h2A3AOkgxIYe9E3Yeoi4BIpc4mU0ZxNw3AziOL9q1QvJyXbJtPsv-DRyWNkNeou3qkTT2EfuyR0eOuUnMviJWMl_nkDcX53MKqYqW0CaUOFwStnKRMlXKcgkhRSlGPwuWFYc9P5J-x-r5zT6e7jC3ugsiqneIxwWkVOiHdzwCUtO_eHtulYQcnhTiBumcgdpujDTPHIcFqWzoios9inhZ04mdWJtUUAjmKUSB-RjzaHUQ0oMK0SoyLtR4e_1we2GPqKOFMmXC2k01wZiqP3FOVcFlfs--oNlfPns8mfzw7nktcXOhvHo8-qG2g0CyVzUDILvVDewrwWQXiTF1f4vCZB9DaNrjPmFQpmEs97NB57tKDa4HdbIDZyeIji49TTfurqDuWdpvk-Za9ocof1Tqf9LWtBObfJXjHOn7vVJTocuaSmXweudehjcc0zRLZzzGTv_EtSvLTyfWJ-TN5qLI_8_-UlY3nhn8sbNX5t3queaZv3-O0cfvsGbzsap60o6my6cmsUP9jvZCWftFe3qGnMhFfaCe0st_pO_qbY1ApoOYVHd-FTSVSd9N-8i_DI9Z3qSKYMQzin9UEe7tZHo5gw_3TIow1_Z9x6NlsGnntpOV2UhiVpWIAaxS7UwLu5G94zq_Pd0CD7FE0e7WKcsSqw2_bA_DRuU7qtjWlcCrkN-ImZus2DQp4ROXJ-uf4sGyX_A4VB5FgxDhqRYxSvySrdEESOb5hvCmPWuDnHNkbdWDb2aL_ZnT_juOu3GhRgqsDlMeCS2f01CIOZ1i1orIBTA6WNv6mZxo2SOYdzd4aKgrnTUSlBC0Q2Bis4MW1A4a_744FR7lzSHQRIgAspLqCMHf0mvz59fv7RR_5zhZvWYFwooAYGAKaV5fOP937-P56-d6NPVOvgvUH8yjjHhTw3nDLhTpLeOA4nyu1WuDsj4ltN3iiSQ2X3X-8Ic6esYcZYvrfb1kwfbIvWcry1Ihu_tdrY5mJq8AdZx5DDlQRK7PUxUnZRWQX4u1Dg7Cm9ZppJsafa_LvBBRVCGsug4CwvUPqTdKuhtHIaagwogTl7ARfyYDZ9ul7il9m8ZdwwEbS91KzoxWbX_mJbYMbifrtj3f62xa8_wJgMbd-RrVc3ui7KXVymcUoXsIs2IYlWUbQmi3pXkjXZrNYQw2oLWxqSqiBVURa0WqcRIemC7UhI4oiEabROUkKCfFuSMI2TakvDsMpXaBXCmTIe2HoPpDotXMXsNqs0Thec5sC1e5lBiC0RFD-4sxgiBCWHhdq5NpG3J41WIWfa6IHIMMPdi5BuYj83OcyVm38dIkWXuotW8d2H25Qvd0SOTv__BQAA___yBRqi">