gistfile1.txt

-:17:12: error: 'spv.Store' op mismatch in result type and pointer type
-:17:12: note: see current operation: "spv.Store"(%9, %0) : (!spv.ptr<i64, StorageBuffer>, i32) -> ()
-:10:1: error: failed to run translation of source executable to target executable for backend #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>
-:10:1: note: see current operation: "hal.executable.variant"() ({
  "hal.executable.entry_point"() ({
  ^bb0(%arg0: !hal.device, %arg1: index, %arg2: index, %arg3: index):
    %0 = "arith.constant"() {value = 1 : index} : () -> index
    "hal.return"(%0, %0, %0) : (index, index, index) -> ()
  }) {layout = #hal.executable.layout<push_constants = 0, sets = [#hal.descriptor_set.layout<0, bindings = [#hal.descriptor_set.binding<0, storage_buffer>]>]>, ordinal = 0 : index, sym_name = "__builtin_splat_i64", translation_info = #iree_codegen.translation_info<SPIRVDistribute>, workgroup_size = [1 : index, 1 : index, 1 : index]} : () -> ()
  "builtin.module"() ({
    "spv.module"() ({
      "spv.GlobalVariable"() {binding = 0 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_0_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
      "spv.func"() ({
        %0 = "spv.Constant"() {value = 12 : i32} : () -> i32
        %1 = "spv.Constant"() {value = 0 : i32} : () -> i32
        %2 = "spv.Constant"() {value = 1 : i32} : () -> i32
        %3 = "spv.mlir.addressof"() {variable = @__resource_var_0_0_} : () -> !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>
        %4 = "spv.Constant"() {value = 0 : i32} : () -> i32
        %5 = "spv.Constant"() {value = 0 : i32} : () -> i32
        %6 = "spv.Constant"() {value = 1 : i32} : () -> i32
        %7 = "spv.IMul"(%6, %1) : (i32, i32) -> i32
        %8 = "spv.IAdd"(%5, %7) : (i32, i32) -> i32
        %9 = "spv.AccessChain"(%3, %4, %8) : (!spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>, i32, i32) -> !spv.ptr<i64, StorageBuffer>
        "spv.Store"(%9, %0) : (!spv.ptr<i64, StorageBuffer>, i32) -> ()
        "spv.Return"() : () -> ()
      }) {function_control = 0 : i32, function_type = () -> (), spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, sym_name = "__builtin_splat_i64"} : () -> ()
    }) {addressing_model = 0 : i32, memory_model = 1 : i32} : () -> ()
  }) {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}>} : () -> ()
  "hal.executable.variant_end"() : () -> ()
}) {sym_name = "vulkan_spirv_fb", target = #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>} : () -> ()
-:10:1: error: failed to serialize executables
-:10:1: note: see current operation: "hal.executable"() ({
  "hal.executable.variant"() ({
    "hal.executable.entry_point"() ({
    ^bb0(%arg0: !hal.device, %arg1: index, %arg2: index, %arg3: index):
      %0 = "arith.constant"() {value = 1 : index} : () -> index
      "hal.return"(%0, %0, %0) : (index, index, index) -> ()
    }) {layout = #hal.executable.layout<push_constants = 0, sets = [#hal.descriptor_set.layout<0, bindings = [#hal.descriptor_set.binding<0, storage_buffer>]>]>, ordinal = 0 : index, sym_name = "__builtin_splat_i64", translation_info = #iree_codegen.translation_info<SPIRVDistribute>, workgroup_size = [1 : index, 1 : index, 1 : index]} : () -> ()
    "builtin.module"() ({
      "spv.module"() ({
        "spv.GlobalVariable"() {binding = 0 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_0_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
        "spv.func"() ({
          %0 = "spv.Constant"() {value = 12 : i32} : () -> i32
          %1 = "spv.Constant"() {value = 0 : i32} : () -> i32
          %2 = "spv.Constant"() {value = 1 : i32} : () -> i32
          %3 = "spv.mlir.addressof"() {variable = @__resource_var_0_0_} : () -> !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>
          %4 = "spv.Constant"() {value = 0 : i32} : () -> i32
          %5 = "spv.Constant"() {value = 0 : i32} : () -> i32
          %6 = "spv.Constant"() {value = 1 : i32} : () -> i32
          %7 = "spv.IMul"(%6, %1) : (i32, i32) -> i32
          %8 = "spv.IAdd"(%5, %7) : (i32, i32) -> i32
          %9 = "spv.AccessChain"(%3, %4, %8) : (!spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>, i32, i32) -> !spv.ptr<i64, StorageBuffer>
          "spv.Store"(%9, %0) : (!spv.ptr<i64, StorageBuffer>, i32) -> ()
          "spv.Return"() : () -> ()
        }) {function_control = 0 : i32, function_type = () -> (), spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, sym_name = "__builtin_splat_i64"} : () -> ()
      }) {addressing_model = 0 : i32, memory_model = 1 : i32} : () -> ()
    }) {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}>} : () -> ()
    "hal.executable.variant_end"() : () -> ()
  }) {sym_name = "vulkan_spirv_fb", target = #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>} : () -> ()
  "hal.executable_end"() : () -> ()
}) {sym_name = "__builtin_splat_i64", sym_visibility = "private"} : () -> ()
/tmp/iree-temps/core-input.mlir:241:11: error: 'spv.Store' op mismatch in result type and pointer type
    %75 = linalg.generic {indexing_maps = [#map5, #map5, #map5], iterator_types = []} ins(%74, %cst_9 : tensor<i64>, tensor<i64>) outs(%73 : tensor<i64>) {
          ^
/tmp/iree-temps/core-input.mlir:23:3: note: called from
  func.func @forward(%arg0: tensor<1x12xi64>) -> tensor<1x2xf32> {
  ^
/tmp/iree-temps/core-input.mlir:241:11: note: see current operation: "spv.Store"(%11, %3) : (!spv.ptr<i64, StorageBuffer>, i32) -> ()
    %75 = linalg.generic {indexing_maps = [#map5, #map5, #map5], iterator_types = []} ins(%74, %cst_9 : tensor<i64>, tensor<i64>) outs(%73 : tensor<i64>) {
          ^
/tmp/iree-temps/core-input.mlir:241:11: error: failed to run translation of source executable to target executable for backend #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>
    %75 = linalg.generic {indexing_maps = [#map5, #map5, #map5], iterator_types = []} ins(%74, %cst_9 : tensor<i64>, tensor<i64>) outs(%73 : tensor<i64>) {
          ^
/tmp/iree-temps/core-input.mlir:23:3: note: called from
  func.func @forward(%arg0: tensor<1x12xi64>) -> tensor<1x2xf32> {
  ^
/tmp/iree-temps/core-input.mlir:241:11: note: see current operation: "hal.executable.variant"() ({
  "hal.executable.entry_point"() ({
  ^bb0(%arg0: !hal.device, %arg1: index, %arg2: index, %arg3: index):
    %0 = "arith.constant"() {value = 1 : index} : () -> index
    "hal.return"(%0, %0, %0) : (index, index, index) -> ()
  }) {layout = #hal.executable.layout<push_constants = 0, sets = [#hal.descriptor_set.layout<0, bindings = [#hal.descriptor_set.binding<0, storage_buffer>, #hal.descriptor_set.binding<1, storage_buffer>]>]>, ordinal = 0 : index, sym_name = "forward_dispatch_0", translation_info = #iree_codegen.translation_info<SPIRVDistribute>, workgroup_size = [1 : index, 1 : index, 1 : index]} : () -> ()
  "builtin.module"() ({
    "spv.module"() ({
      "spv.GlobalVariable"() {binding = 1 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_1_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
      "spv.func"() ({
        %0 = "spv.Constant"() {value = 0 : i32} : () -> i32
        %1 = "spv.Constant"() {value = 8 : i32} : () -> i32
        %2 = "spv.Constant"() {value = 64 : i32} : () -> i32
        %3 = "spv.Constant"() {value = 12 : i32} : () -> i32
        %4 = "spv.Constant"() {value = 1 : i32} : () -> i32
        %5 = "spv.mlir.addressof"() {variable = @__resource_var_0_1_} : () -> !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>
        %6 = "spv.Constant"() {value = 0 : i32} : () -> i32
        %7 = "spv.Constant"() {value = 0 : i32} : () -> i32
        %8 = "spv.Constant"() {value = 1 : i32} : () -> i32
        %9 = "spv.IMul"(%8, %1) : (i32, i32) -> i32
        %10 = "spv.IAdd"(%7, %9) : (i32, i32) -> i32
        %11 = "spv.AccessChain"(%5, %6, %10) : (!spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>, i32, i32) -> !spv.ptr<i64, StorageBuffer>
        "spv.Store"(%11, %3) : (!spv.ptr<i64, StorageBuffer>, i32) -> ()
        "spv.Return"() : () -> ()
      }) {function_control = 0 : i32, function_type = () -> (), spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, sym_name = "forward_dispatch_0"} : () -> ()
    }) {addressing_model = 0 : i32, memory_model = 1 : i32} : () -> ()
  }) {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}>} : () -> ()
  "hal.executable.variant_end"() : () -> ()
}) {sym_name = "vulkan_spirv_fb", target = #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>} : () -> ()
    %75 = linalg.generic {indexing_maps = [#map5, #map5, #map5], iterator_types = []} ins(%74, %cst_9 : tensor<i64>, tensor<i64>) outs(%73 : tensor<i64>) {
          ^
/tmp/iree-temps/core-input.mlir:241:11: error: failed to serialize executables
    %75 = linalg.generic {indexing_maps = [#map5, #map5, #map5], iterator_types = []} ins(%74, %cst_9 : tensor<i64>, tensor<i64>) outs(%73 : tensor<i64>) {
          ^
/tmp/iree-temps/core-input.mlir:23:3: note: called from
  func.func @forward(%arg0: tensor<1x12xi64>) -> tensor<1x2xf32> {
  ^
/tmp/iree-temps/core-input.mlir:241:11: note: see current operation: "hal.executable"() ({
  "hal.executable.variant"() ({
    "hal.executable.entry_point"() ({
    ^bb0(%arg0: !hal.device, %arg1: index, %arg2: index, %arg3: index):
      %0 = "arith.constant"() {value = 1 : index} : () -> index
      "hal.return"(%0, %0, %0) : (index, index, index) -> ()
    }) {layout = #hal.executable.layout<push_constants = 0, sets = [#hal.descriptor_set.layout<0, bindings = [#hal.descriptor_set.binding<0, storage_buffer>, #hal.descriptor_set.binding<1, storage_buffer>]>]>, ordinal = 0 : index, sym_name = "forward_dispatch_0", translation_info = #iree_codegen.translation_info<SPIRVDistribute>, workgroup_size = [1 : index, 1 : index, 1 : index]} : () -> ()
    "builtin.module"() ({
      "spv.module"() ({
        "spv.GlobalVariable"() {binding = 1 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_1_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
        "spv.func"() ({
          %0 = "spv.Constant"() {value = 0 : i32} : () -> i32
          %1 = "spv.Constant"() {value = 8 : i32} : () -> i32
          %2 = "spv.Constant"() {value = 64 : i32} : () -> i32
          %3 = "spv.Constant"() {value = 12 : i32} : () -> i32
          %4 = "spv.Constant"() {value = 1 : i32} : () -> i32
          %5 = "spv.mlir.addressof"() {variable = @__resource_var_0_1_} : () -> !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>
          %6 = "spv.Constant"() {value = 0 : i32} : () -> i32
          %7 = "spv.Constant"() {value = 0 : i32} : () -> i32
          %8 = "spv.Constant"() {value = 1 : i32} : () -> i32
          %9 = "spv.IMul"(%8, %1) : (i32, i32) -> i32
          %10 = "spv.IAdd"(%7, %9) : (i32, i32) -> i32
          %11 = "spv.AccessChain"(%5, %6, %10) : (!spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>, i32, i32) -> !spv.ptr<i64, StorageBuffer>
          "spv.Store"(%11, %3) : (!spv.ptr<i64, StorageBuffer>, i32) -> ()
          "spv.Return"() : () -> ()
        }) {function_control = 0 : i32, function_type = () -> (), spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, sym_name = "forward_dispatch_0"} : () -> ()
      }) {addressing_model = 0 : i32, memory_model = 1 : i32} : () -> ()
    }) {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}>} : () -> ()
    "hal.executable.variant_end"() : () -> ()
  }) {sym_name = "vulkan_spirv_fb", target = #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>} : () -> ()
  "hal.executable_end"() : () -> ()
}) {sym_name = "forward_dispatch_0", sym_visibility = "private"} : () -> ()
    %75 = linalg.generic {indexing_maps = [#map5, #map5, #map5], iterator_types = []} ins(%74, %cst_9 : tensor<i64>, tensor<i64>) outs(%73 : tensor<i64>) {
          ^
/tmp/iree-temps/core-input.mlir:261:11: error: failed to legalize operation 'memref.load'
    %90 = tensor.extract_slice %78[0, %29] [1, %89] [1, 1] : tensor<1x512xi64> to tensor<1x?xi64>
          ^
/tmp/iree-temps/core-input.mlir:23:3: note: called from
  func.func @forward(%arg0: tensor<1x12xi64>) -> tensor<1x2xf32> {
  ^
/tmp/iree-temps/core-input.mlir:261:11: note: see current operation: %64 = "memref.load"(%25, %63) : (memref<?xi64>, index) -> i64
    %90 = tensor.extract_slice %78[0, %29] [1, %89] [1, 1] : tensor<1x512xi64> to tensor<1x?xi64>
          ^
/tmp/iree-temps/core-input.mlir:261:11: error: failed to run translation of source executable to target executable for backend #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>
    %90 = tensor.extract_slice %78[0, %29] [1, %89] [1, 1] : tensor<1x512xi64> to tensor<1x?xi64>
          ^
/tmp/iree-temps/core-input.mlir:23:3: note: called from
  func.func @forward(%arg0: tensor<1x12xi64>) -> tensor<1x2xf32> {
  ^
/tmp/iree-temps/core-input.mlir:261:11: note: see current operation: "hal.executable.variant"() ({
  "hal.executable.entry_point"() ({
  ^bb0(%arg0: !hal.device, %arg1: index, %arg2: index, %arg3: index):
    %0 = "arith.constant"() {value = 1 : index} : () -> index
    %1 = "affine.apply"(%arg1) {map = affine_map<()[s0] -> (s0 ceildiv 4)>} : (index) -> index
    "hal.return"(%1, %0, %0) : (index, index, index) -> ()
  }) {layout = #hal.executable.layout<push_constants = 1, sets = [#hal.descriptor_set.layout<0, bindings = [#hal.descriptor_set.binding<0, storage_buffer>, #hal.descriptor_set.binding<1, storage_buffer>]>]>, ordinal = 0 : index, sym_name = "forward_dispatch_1", translation_info = #iree_codegen.translation_info<SPIRVDistribute workload_per_wg = [4]>, workgroup_size = [4 : index, 1 : index, 1 : index]} : () -> ()
  "builtin.module"() ({
    "spv.GlobalVariable"() {binding = 0 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_0_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
    "spv.GlobalVariable"() {binding = 1 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_1_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
    "func.func"() ({
      %0 = "arith.constant"() {value = 13840 : index} : () -> index
      %1 = "arith.constant"() {value = 11356992 : index} : () -> index
      %2 = "arith.constant"() {value = -1 : index} : () -> index
      %3 = "arith.constant"() {value = 4 : index} : () -> index
      %4 = "arith.constant"() {value = 0 : index} : () -> index
      %5 = "arith.constant"() {value = 512 : index} : () -> index
      %6 = "arith.constant"() {value = 90855936 : index} : () -> index
      %7 = "arith.constant"() {value = 110720 : index} : () -> index
      %8 = "hal.interface.constant.load"() {index = 0 : index} : () -> i32
      %9 = "arith.index_cast"(%8) : (i32) -> index
      %10 = "hal.interface.binding.subspan"(%6, %5) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
      %11 = "hal.interface.binding.subspan"(%4, %5) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
      %12 = "hal.interface.binding.subspan"(%7, %9) {alignment = 64 : index, binding = 1 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
      %13 = "hal.interface.binding.subspan"(%4, %9) {alignment = 64 : index, binding = 1 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
      %14 = "hal.interface.workgroup.id"() {dimension = 0 : index} : () -> index
      %15 = "hal.interface.workgroup.count"() {dimension = 0 : index} : () -> index
      %16 = "arith.muli"(%14, %3) : (index, index) -> index
      %17 = "arith.muli"(%15, %3) : (index, index) -> index
      %18 = "gpu.thread_id"() {dimension = #gpu<"dim x">} : () -> index
      %19 = "gpu.block_dim"() {dimension = #gpu<"dim x">} : () -> index
      "scf.for"(%16, %9, %17) ({
      ^bb0(%arg0: index):
        %20 = "arith.muli"(%arg0, %2) : (index, index) -> index
        %21 = "arith.addi"(%20, %9) : (index, index) -> index
        %22 = "arith.cmpi"(%21, %3) {predicate = 2 : i64} : (index, index) -> i1
        %23 = "arith.select"(%22, %21, %3) : (i1, index, index) -> index
        "scf.for"(%18, %23, %19) ({
        ^bb0(%arg1: index):
          %24 = "arith.addi"(%arg1, %arg0) : (index, index) -> index
          %25 = "arith.addi"(%24, %1) : (index, index) -> index
          %26 = "memref.load"(%11, %25) : (memref<?xi64>, index) -> i64
          %27 = "arith.addi"(%24, %0) : (index, index) -> index
          "memref.store"(%26, %13, %27) : (i64, memref<?xi64>, index) -> ()
          "scf.yield"() : () -> ()
        }) : (index, index, index) -> ()
        "scf.yield"() : () -> ()
      }) : (index, index, index) -> ()
      "func.return"() : () -> ()
    }) {function_type = () -> (), spv.entry_point_abi = {local_size = dense<[4, 1, 1]> : vector<3xi32>}, sym_name = "forward_dispatch_1"} : () -> ()
  }) {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}>} : () -> ()
  "hal.executable.variant_end"() : () -> ()
}) {sym_name = "vulkan_spirv_fb", target = #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>} : () -> ()
    %90 = tensor.extract_slice %78[0, %29] [1, %89] [1, 1] : tensor<1x512xi64> to tensor<1x?xi64>
          ^
/tmp/iree-temps/core-input.mlir:261:11: error: failed to serialize executables
    %90 = tensor.extract_slice %78[0, %29] [1, %89] [1, 1] : tensor<1x512xi64> to tensor<1x?xi64>
          ^
/tmp/iree-temps/core-input.mlir:23:3: note: called from
  func.func @forward(%arg0: tensor<1x12xi64>) -> tensor<1x2xf32> {
  ^
/tmp/iree-temps/core-input.mlir:261:11: note: see current operation: "hal.executable"() ({
  "hal.executable.variant"() ({
    "hal.executable.entry_point"() ({
    ^bb0(%arg0: !hal.device, %arg1: index, %arg2: index, %arg3: index):
      %0 = "arith.constant"() {value = 1 : index} : () -> index
      %1 = "affine.apply"(%arg1) {map = affine_map<()[s0] -> (s0 ceildiv 4)>} : (index) -> index
      "hal.return"(%1, %0, %0) : (index, index, index) -> ()
    }) {layout = #hal.executable.layout<push_constants = 1, sets = [#hal.descriptor_set.layout<0, bindings = [#hal.descriptor_set.binding<0, storage_buffer>, #hal.descriptor_set.binding<1, storage_buffer>]>]>, ordinal = 0 : index, sym_name = "forward_dispatch_1", translation_info = #iree_codegen.translation_info<SPIRVDistribute workload_per_wg = [4]>, workgroup_size = [4 : index, 1 : index, 1 : index]} : () -> ()
    "builtin.module"() ({
      "spv.GlobalVariable"() {binding = 0 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_0_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
      "spv.GlobalVariable"() {binding = 1 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_1_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
      "func.func"() ({
        %0 = "arith.constant"() {value = 13840 : index} : () -> index
        %1 = "arith.constant"() {value = 11356992 : index} : () -> index
        %2 = "arith.constant"() {value = -1 : index} : () -> index
        %3 = "arith.constant"() {value = 4 : index} : () -> index
        %4 = "arith.constant"() {value = 0 : index} : () -> index
        %5 = "arith.constant"() {value = 512 : index} : () -> index
        %6 = "arith.constant"() {value = 90855936 : index} : () -> index
        %7 = "arith.constant"() {value = 110720 : index} : () -> index
        %8 = "hal.interface.constant.load"() {index = 0 : index} : () -> i32
        %9 = "arith.index_cast"(%8) : (i32) -> index
        %10 = "hal.interface.binding.subspan"(%6, %5) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
        %11 = "hal.interface.binding.subspan"(%4, %5) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
        %12 = "hal.interface.binding.subspan"(%7, %9) {alignment = 64 : index, binding = 1 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
        %13 = "hal.interface.binding.subspan"(%4, %9) {alignment = 64 : index, binding = 1 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
        %14 = "hal.interface.workgroup.id"() {dimension = 0 : index} : () -> index
        %15 = "hal.interface.workgroup.count"() {dimension = 0 : index} : () -> index
        %16 = "arith.muli"(%14, %3) : (index, index) -> index
        %17 = "arith.muli"(%15, %3) : (index, index) -> index
        %18 = "gpu.thread_id"() {dimension = #gpu<"dim x">} : () -> index
        %19 = "gpu.block_dim"() {dimension = #gpu<"dim x">} : () -> index
        "scf.for"(%16, %9, %17) ({
        ^bb0(%arg0: index):
          %20 = "arith.muli"(%arg0, %2) : (index, index) -> index
          %21 = "arith.addi"(%20, %9) : (index, index) -> index
          %22 = "arith.cmpi"(%21, %3) {predicate = 2 : i64} : (index, index) -> i1
          %23 = "arith.select"(%22, %21, %3) : (i1, index, index) -> index
          "scf.for"(%18, %23, %19) ({
          ^bb0(%arg1: index):
            %24 = "arith.addi"(%arg1, %arg0) : (index, index) -> index
            %25 = "arith.addi"(%24, %1) : (index, index) -> index
            %26 = "memref.load"(%11, %25) : (memref<?xi64>, index) -> i64
            %27 = "arith.addi"(%24, %0) : (index, index) -> index
            "memref.store"(%26, %13, %27) : (i64, memref<?xi64>, index) -> ()
            "scf.yield"() : () -> ()
          }) : (index, index, index) -> ()
          "scf.yield"() : () -> ()
        }) : (index, index, index) -> ()
        "func.return"() : () -> ()
      }) {function_type = () -> (), spv.entry_point_abi = {local_size = dense<[4, 1, 1]> : vector<3xi32>}, sym_name = "forward_dispatch_1"} : () -> ()
    }) {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}>} : () -> ()
    "hal.executable.variant_end"() : () -> ()
  }) {sym_name = "vulkan_spirv_fb", target = #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>} : () -> ()
  "hal.executable_end"() : () -> ()
}) {sym_name = "forward_dispatch_1", sym_visibility = "private"} : () -> ()
    %90 = tensor.extract_slice %78[0, %29] [1, %89] [1, 1] : tensor<1x512xi64> to tensor<1x?xi64>
          ^
/tmp/iree-temps/core-input.mlir:304:11: error: failed to legalize operation 'memref.load'
    %97 = linalg.generic {indexing_maps = [#map8, #map8, #map7], iterator_types = ["parallel", "parallel", "parallel"]} ins(%95, %96 : tensor<1x12x384xf32>, tensor<1x12x384xf32>) outs(%92 : tensor<1x12x384xf32>) {
          ^
/tmp/iree-temps/core-input.mlir:23:3: note: called from
  func.func @forward(%arg0: tensor<1x12xi64>) -> tensor<1x2xf32> {
  ^
/tmp/iree-temps/core-input.mlir:304:11: note: see current operation: %71 = "memref.load"(%49, %62) : (memref<?xi64>, index) -> i64
    %97 = linalg.generic {indexing_maps = [#map8, #map8, #map7], iterator_types = ["parallel", "parallel", "parallel"]} ins(%95, %96 : tensor<1x12x384xf32>, tensor<1x12x384xf32>) outs(%92 : tensor<1x12x384xf32>) {
          ^
/tmp/iree-temps/core-input.mlir:304:11: error: failed to run translation of source executable to target executable for backend #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>
    %97 = linalg.generic {indexing_maps = [#map8, #map8, #map7], iterator_types = ["parallel", "parallel", "parallel"]} ins(%95, %96 : tensor<1x12x384xf32>, tensor<1x12x384xf32>) outs(%92 : tensor<1x12x384xf32>) {
          ^
/tmp/iree-temps/core-input.mlir:23:3: note: called from
  func.func @forward(%arg0: tensor<1x12xi64>) -> tensor<1x2xf32> {
  ^
/tmp/iree-temps/core-input.mlir:304:11: note: see current operation: "hal.executable.variant"() ({
  "hal.executable.entry_point"() ({
  ^bb0(%arg0: !hal.device, %arg1: index, %arg2: index, %arg3: index):
    %0 = "arith.constant"() {value = 1 : index} : () -> index
    %1 = "affine.apply"(%arg1) {map = affine_map<()[s0] -> (s0 ceildiv 4)>} : (index) -> index
    "hal.return"(%1, %arg2, %0) : (index, index, index) -> ()
  }) {layout = #hal.executable.layout<push_constants = 1, sets = [#hal.descriptor_set.layout<0, bindings = [#hal.descriptor_set.binding<0, storage_buffer>, #hal.descriptor_set.binding<1, storage_buffer>, #hal.descriptor_set.binding<2, storage_buffer>, #hal.descriptor_set.binding<3, storage_buffer>]>]>, ordinal = 0 : index, sym_name = "forward_dispatch_2", translation_info = #iree_codegen.translation_info<SPIRVDistribute workload_per_wg = [4, 1]>, workgroup_size = [4 : index, 1 : index, 1 : index]} : () -> ()
  "builtin.module"() ({
    "spv.GlobalVariable"() {binding = 0 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_0_", type = !spv.ptr<!spv.struct<(!spv.rtarray<f32, stride=4> [0])>, StorageBuffer>} : () -> ()
    "spv.GlobalVariable"() {binding = 1 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_1_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
    "spv.GlobalVariable"() {binding = 3 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_3_", type = !spv.ptr<!spv.struct<(!spv.rtarray<f32, stride=4> [0])>, StorageBuffer>} : () -> ()
    "spv.GlobalVariable"() {binding = 2 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_2_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
    "func.func"() ({
      %0 = "arith.constant"() {value = 4 : index} : () -> index
      %1 = "arith.constant"() {value = 384 : index} : () -> index
      %2 = "arith.constant"() {value = 10796160 : index} : () -> index
      %3 = "arith.constant"() {value = 10992768 : index} : () -> index
      %4 = "arith.constant"() {value = 10993536 : index} : () -> index
      %5 = "arith.constant"() {value = 13840 : index} : () -> index
      %6 = "arith.constant"() {value = 0 : index} : () -> index
      %7 = "arith.constant"() {value = 4608 : index} : () -> index
      %8 = "arith.constant"() {value = 12 : index} : () -> index
      %9 = "arith.constant"() {value = 196608 : index} : () -> index
      %10 = "arith.constant"() {value = 768 : index} : () -> index
      %11 = "arith.constant"() {value = 11720448 : index} : () -> index
      %12 = "arith.constant"() {value = 43974144 : index} : () -> index
      %13 = "arith.constant"() {value = 43971072 : index} : () -> index
      %14 = "arith.constant"() {value = 43184640 : index} : () -> index
      %15 = "arith.constant"() {value = 110720 : index} : () -> index
      %16 = "hal.interface.constant.load"() {index = 0 : index} : () -> i32
      %17 = "arith.index_cast"(%16) : (i32) -> index
      %18 = "hal.interface.binding.subspan"(%12, %11) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
      %19 = "hal.interface.binding.subspan"(%6, %11) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
      %20 = "hal.interface.binding.subspan"(%13, %10) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
      %21 = "hal.interface.binding.subspan"(%6, %10) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
      %22 = "hal.interface.binding.subspan"(%14, %9) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
      %23 = "hal.interface.binding.subspan"(%6, %9) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
      %24 = "hal.interface.binding.subspan"(%6, %8) {alignment = 64 : index, binding = 1 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
      %25 = "hal.interface.binding.subspan"(%6, %7) {alignment = 64 : index, binding = 3 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
      %26 = "hal.interface.binding.subspan"(%15, %17) {alignment = 64 : index, binding = 2 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
      %27 = "hal.interface.binding.subspan"(%6, %17) {alignment = 64 : index, binding = 2 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
      %28 = "hal.interface.workgroup.id"() {dimension = 0 : index} : () -> index
      %29 = "hal.interface.workgroup.id"() {dimension = 1 : index} : () -> index
      %30 = "gpu.thread_id"() {dimension = #gpu<"dim x">} : () -> index
      %31 = "gpu.thread_id"() {dimension = #gpu<"dim y">} : () -> index
      %32 = "memref.load"(%24, %29) : (memref<?xi64>, index) -> i64
      %33 = "arith.addi"(%29, %5) : (index, index) -> index
      %34 = "memref.load"(%27, %33) : (memref<?xi64>, index) -> i64
      %35 = "arith.index_cast"(%32) : (i64) -> index
      %36 = "arith.muli"(%28, %0) : (index, index) -> index
      %37 = "arith.addi"(%36, %30) : (index, index) -> index
      %38 = "arith.muli"(%35, %1) : (index, index) -> index
      %39 = "arith.addi"(%37, %38) : (index, index) -> index
      %40 = "arith.addi"(%39, %4) : (index, index) -> index
      %41 = "memref.load"(%19, %40) : (memref<?xf32>, index) -> f32
      %42 = "arith.addi"(%37, %3) : (index, index) -> index
      %43 = "memref.load"(%21, %42) : (memref<?xf32>, index) -> f32
      %44 = "arith.index_cast"(%34) : (i64) -> index
      %45 = "arith.muli"(%44, %1) : (index, index) -> index
      %46 = "arith.addi"(%37, %45) : (index, index) -> index
      %47 = "arith.addi"(%46, %2) : (index, index) -> index
      %48 = "memref.load"(%23, %47) : (memref<?xf32>, index) -> f32
      %49 = "arith.addf"(%41, %43) : (f32, f32) -> f32
      %50 = "arith.addf"(%49, %48) : (f32, f32) -> f32
      %51 = "arith.muli"(%29, %1) : (index, index) -> index
      %52 = "arith.muli"(%31, %1) : (index, index) -> index
      %53 = "arith.addi"(%51, %52) : (index, index) -> index
      %54 = "arith.addi"(%53, %36) : (index, index) -> index
      %55 = "arith.addi"(%54, %30) : (index, index) -> index
      "memref.store"(%50, %25, %55) : (f32, memref<?xf32>, index) -> ()
      "func.return"() : () -> ()
    }) {function_type = () -> (), spv.entry_point_abi = {local_size = dense<[4, 1, 1]> : vector<3xi32>}, sym_name = "forward_dispatch_2"} : () -> ()
  }) {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}>} : () -> ()
  "hal.executable.variant_end"() : () -> ()
}) {sym_name = "vulkan_spirv_fb", target = #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>} : () -> ()
    %97 = linalg.generic {indexing_maps = [#map8, #map8, #map7], iterator_types = ["parallel", "parallel", "parallel"]} ins(%95, %96 : tensor<1x12x384xf32>, tensor<1x12x384xf32>) outs(%92 : tensor<1x12x384xf32>) {
          ^
/tmp/iree-temps/core-input.mlir:304:11: error: failed to serialize executables
    %97 = linalg.generic {indexing_maps = [#map8, #map8, #map7], iterator_types = ["parallel", "parallel", "parallel"]} ins(%95, %96 : tensor<1x12x384xf32>, tensor<1x12x384xf32>) outs(%92 : tensor<1x12x384xf32>) {
          ^
/tmp/iree-temps/core-input.mlir:23:3: note: called from
  func.func @forward(%arg0: tensor<1x12xi64>) -> tensor<1x2xf32> {
  ^
/tmp/iree-temps/core-input.mlir:304:11: note: see current operation: "hal.executable"() ({
  "hal.executable.variant"() ({
    "hal.executable.entry_point"() ({
    ^bb0(%arg0: !hal.device, %arg1: index, %arg2: index, %arg3: index):
      %0 = "arith.constant"() {value = 1 : index} : () -> index
      %1 = "affine.apply"(%arg1) {map = affine_map<()[s0] -> (s0 ceildiv 4)>} : (index) -> index
      "hal.return"(%1, %arg2, %0) : (index, index, index) -> ()
    }) {layout = #hal.executable.layout<push_constants = 1, sets = [#hal.descriptor_set.layout<0, bindings = [#hal.descriptor_set.binding<0, storage_buffer>, #hal.descriptor_set.binding<1, storage_buffer>, #hal.descriptor_set.binding<2, storage_buffer>, #hal.descriptor_set.binding<3, storage_buffer>]>]>, ordinal = 0 : index, sym_name = "forward_dispatch_2", translation_info = #iree_codegen.translation_info<SPIRVDistribute workload_per_wg = [4, 1]>, workgroup_size = [4 : index, 1 : index, 1 : index]} : () -> ()
    "builtin.module"() ({
      "spv.GlobalVariable"() {binding = 0 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_0_", type = !spv.ptr<!spv.struct<(!spv.rtarray<f32, stride=4> [0])>, StorageBuffer>} : () -> ()
      "spv.GlobalVariable"() {binding = 1 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_1_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
      "spv.GlobalVariable"() {binding = 3 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_3_", type = !spv.ptr<!spv.struct<(!spv.rtarray<f32, stride=4> [0])>, StorageBuffer>} : () -> ()
      "spv.GlobalVariable"() {binding = 2 : i32, descriptor_set = 0 : i32, sym_name = "__resource_var_0_2_", type = !spv.ptr<!spv.struct<(!spv.rtarray<i64, stride=8> [0])>, StorageBuffer>} : () -> ()
      "func.func"() ({
        %0 = "arith.constant"() {value = 4 : index} : () -> index
        %1 = "arith.constant"() {value = 384 : index} : () -> index
        %2 = "arith.constant"() {value = 10796160 : index} : () -> index
        %3 = "arith.constant"() {value = 10992768 : index} : () -> index
        %4 = "arith.constant"() {value = 10993536 : index} : () -> index
        %5 = "arith.constant"() {value = 13840 : index} : () -> index
        %6 = "arith.constant"() {value = 0 : index} : () -> index
        %7 = "arith.constant"() {value = 4608 : index} : () -> index
        %8 = "arith.constant"() {value = 12 : index} : () -> index
        %9 = "arith.constant"() {value = 196608 : index} : () -> index
        %10 = "arith.constant"() {value = 768 : index} : () -> index
        %11 = "arith.constant"() {value = 11720448 : index} : () -> index
        %12 = "arith.constant"() {value = 43974144 : index} : () -> index
        %13 = "arith.constant"() {value = 43971072 : index} : () -> index
        %14 = "arith.constant"() {value = 43184640 : index} : () -> index
        %15 = "arith.constant"() {value = 110720 : index} : () -> index
        %16 = "hal.interface.constant.load"() {index = 0 : index} : () -> i32
        %17 = "arith.index_cast"(%16) : (i32) -> index
        %18 = "hal.interface.binding.subspan"(%12, %11) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
        %19 = "hal.interface.binding.subspan"(%6, %11) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
        %20 = "hal.interface.binding.subspan"(%13, %10) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
        %21 = "hal.interface.binding.subspan"(%6, %10) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
        %22 = "hal.interface.binding.subspan"(%14, %9) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
        %23 = "hal.interface.binding.subspan"(%6, %9) {alignment = 64 : index, binding = 0 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
        %24 = "hal.interface.binding.subspan"(%6, %8) {alignment = 64 : index, binding = 1 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
        %25 = "hal.interface.binding.subspan"(%6, %7) {alignment = 64 : index, binding = 3 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xf32>
        %26 = "hal.interface.binding.subspan"(%15, %17) {alignment = 64 : index, binding = 2 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
        %27 = "hal.interface.binding.subspan"(%6, %17) {alignment = 64 : index, binding = 2 : index, operand_segment_sizes = dense<1> : vector<2xi32>, set = 0 : index, type = 7 : i32} : (index, index) -> memref<?xi64>
        %28 = "hal.interface.workgroup.id"() {dimension = 0 : index} : () -> index
        %29 = "hal.interface.workgroup.id"() {dimension = 1 : index} : () -> index
        %30 = "gpu.thread_id"() {dimension = #gpu<"dim x">} : () -> index
        %31 = "gpu.thread_id"() {dimension = #gpu<"dim y">} : () -> index
        %32 = "memref.load"(%24, %29) : (memref<?xi64>, index) -> i64
        %33 = "arith.addi"(%29, %5) : (index, index) -> index
        %34 = "memref.load"(%27, %33) : (memref<?xi64>, index) -> i64
        %35 = "arith.index_cast"(%32) : (i64) -> index
        %36 = "arith.muli"(%28, %0) : (index, index) -> index
        %37 = "arith.addi"(%36, %30) : (index, index) -> index
        %38 = "arith.muli"(%35, %1) : (index, index) -> index
        %39 = "arith.addi"(%37, %38) : (index, index) -> index
        %40 = "arith.addi"(%39, %4) : (index, index) -> index
        %41 = "memref.load"(%19, %40) : (memref<?xf32>, index) -> f32
        %42 = "arith.addi"(%37, %3) : (index, index) -> index
        %43 = "memref.load"(%21, %42) : (memref<?xf32>, index) -> f32
        %44 = "arith.index_cast"(%34) : (i64) -> index
        %45 = "arith.muli"(%44, %1) : (index, index) -> index
        %46 = "arith.addi"(%37, %45) : (index, index) -> index
        %47 = "arith.addi"(%46, %2) : (index, index) -> index
        %48 = "memref.load"(%23, %47) : (memref<?xf32>, index) -> f32
        %49 = "arith.addf"(%41, %43) : (f32, f32) -> f32
        %50 = "arith.addf"(%49, %48) : (f32, f32) -> f32
        %51 = "arith.muli"(%29, %1) : (index, index) -> index
        %52 = "arith.muli"(%31, %1) : (index, index) -> index
        %53 = "arith.addi"(%51, %52) : (index, index) -> index
        %54 = "arith.addi"(%53, %36) : (index, index) -> index
        %55 = "arith.addi"(%54, %30) : (index, index) -> index
        "memref.store"(%50, %25, %55) : (f32, memref<?xf32>, index) -> ()
        "func.return"() : () -> ()
      }) {function_type = () -> (), spv.entry_point_abi = {local_size = dense<[4, 1, 1]> : vector<3xi32>}, sym_name = "forward_dispatch_2"} : () -> ()
    }) {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}>} : () -> ()
    "hal.executable.variant_end"() : () -> ()
  }) {sym_name = "vulkan_spirv_fb", target = #hal.executable.target<"vulkan", "vulkan-spirv-fb", {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}>}>} : () -> ()
  "hal.executable_end"() : () -> ()
}) {sym_name = "forward_dispatch_2", sym_visibility = "private"} : () -> ()
    %97 = linalg.generic {indexing_maps = [#map8, #map8, #map7], iterator_types = ["parallel", "parallel", "parallel"]} ins(%95, %96 : tensor<1x12x384xf32>, tensor<1x12x384xf32>) outs(%92 : tensor<1x12x384xf32>) {
          ^
compilation from source to vm failed